diff -Naur CITK_lastest_devel_build/CMakeLists.txt CITK_CudaCanny//CMakeLists.txt
--- CITK_lastest_devel_build/CMakeLists.txt	2010-05-07 11:09:30.000000000 -0300
+++ CITK_CudaCanny//CMakeLists.txt	2011-08-07 19:04:20.085186000 -0300
@@ -47,7 +47,7 @@
 #-----------------------------------------------------------------------------
 # ITK version number.  An even minor number corresponds to releases.
 SET(ITK_VERSION_MAJOR "3")
-SET(ITK_VERSION_MINOR "16")
+SET(ITK_VERSION_MINOR "20")
 SET(ITK_VERSION_PATCH "0")
 
 # Version string should not include patch level.  The major.minor is
@@ -110,7 +110,13 @@
 MARK_AS_ADVANCED(USE_FFTWF)
 OPTION(USE_FFTWF "Use single precision FFTW if found" OFF)
 IF( USE_FFTWD OR USE_FFTWF )
-    FIND_PACKAGE( FFTW )
+  SET(msg "ATTENTION: You have enabled the use of FFTW,")
+  SET(msg "${msg} this library is distributed under a GPL license.")
+  SET(msg "${msg} By enabling this option, the binary of the ITK libraries")
+  SET(msg "${msg} that you are going to build will be covered by a GPL license,")
+  SET(msg "${msg} and so it will be any executable that you link against these libraries.")
+  MESSAGE("${msg}")
+  FIND_PACKAGE( FFTW )
 ENDIF( USE_FFTWD OR USE_FFTWF )
 
 # Deal with SGI SCSL Options
@@ -131,9 +137,6 @@
 ENDIF( ITK_USE_MINC2 )
 
 
-# Put "." in the rpath of all ITK shared libraries.
-LINK_DIRECTORIES(.)
-
 #-----------------------------------------------------------------------------
 # ITK build configuration options.
 OPTION(BUILD_SHARED_LIBS "Build ITK with shared libraries." OFF)
@@ -278,6 +281,10 @@
 OPTION(ITK_USE_DEPRECATED_LEVELSET_INTERPOLATION "Turn on the use of a sparse field level set surface interpolation method that has been deprecated since ITK 3.6. Turn this flag ON to exactly replicate the behavior of ITK<=3.4." OFF)
 MARK_AS_ADVANCED(ITK_USE_DEPRECATED_LEVELSET_INTERPOLATION)
 
+OPTION(ITK_USE_DEPRECATED_FAST_MARCHING "Turn on incorrect handling of TrialPoints in itk::FastMarchingImageFilter result of a bug fixed in ITK 3.18. Turn this flag ON to exactly replicate the behavior of ITK<=3.16." OFF)
+MARK_AS_ADVANCED(ITK_USE_DEPRECATED_FAST_MARCHING)
+
+
 #-----------------------------------------------------------------------------
 # Installation vars.
 # ITK_INSTALL_BIN_DIR          - binary dir (executables)
@@ -289,6 +296,9 @@
 # ITK_INSTALL_NO_RUNTIME       - do not install runtime files
 # ITK_INSTALL_NO_DOCUMENTATION - do not install documentation files
 
+OPTION(ITK_USE_FLAT_DIRECTORY_INSTALL  "When ON, ITK will install all the header files in a single directory. This reduces the number of directories that must be searched for headers in projects based on ITK" OFF)
+MARK_AS_ADVANCED(ITK_USE_FLAT_DIRECTORY_INSTALL)
+
 IF(NOT ITK_INSTALL_BIN_DIR)
   # Solaris: unlike libraries, 64 bit executables should be installed in "bin${sun64}" only
   # if the 32 version is installed in  "bin/sparcv7" or "bin/i86", and a isaexec
@@ -650,75 +660,21 @@
 
 INCLUDE(CheckTypeSize)
 
-CHECK_TYPE_SIZE("long long" ITK_SIZEOF_LONG_LONG)
-CHECK_TYPE_SIZE("__int64"   ITK_SIZEOF___INT64)
 
-IF(ITK_SIZEOF___INT64)
-  IF("ITK_TYPE_SAME_LONG_AND___INT64" MATCHES "^ITK_TYPE_SAME_LONG_AND___INT64$")
-    MESSAGE(STATUS "Checking whether long and __int64 are the same type")
-    TRY_COMPILE(ITK_TYPE_SAME_LONG_AND___INT64
-      ${ITK_BINARY_DIR}/CMakeTmp
-      ${ITK_SOURCE_DIR}/CMake/itkTestCompareTypes.cxx
-      COMPILE_DEFINITIONS
-      -DITK_TEST_COMPARE_TYPE_1=long
-      -DITK_TEST_COMPARE_TYPE_2=__int64
-      OUTPUT_VARIABLE OUTPUT)
-    IF(ITK_TYPE_SAME_LONG_AND___INT64)
-      MESSAGE(STATUS "Checking whether long and __int64 are the same type -- yes")
-      SET(ITK_TYPE_SAME_LONG_AND___INT64 1 CACHE INTERNAL "Whether long and __int64 are the same type")
-    ELSE(ITK_TYPE_SAME_LONG_AND___INT64)
-      MESSAGE(STATUS "Checking whether long and __int64 are the same type -- no")
-      SET(ITK_TYPE_SAME_LONG_AND___INT64 0 CACHE INTERNAL "Whether long and __int64 are the same type")
-      WRITE_FILE(${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeError.log
-        "Determining whether long and __int64 are the same type "
-        "failed with the following output:\n"
-        "${OUTPUT}\n" APPEND)
-    ENDIF(ITK_TYPE_SAME_LONG_AND___INT64)
-  ENDIF("ITK_TYPE_SAME_LONG_AND___INT64" MATCHES "^ITK_TYPE_SAME_LONG_AND___INT64$")
-ENDIF(ITK_SIZEOF___INT64)
-
-IF(ITK_SIZEOF_LONG_LONG)
-  CONFIGURE_FILE(${ITK_SOURCE_DIR}/CMake/itkTestStreamLongLong.cxx.in
-    ${ITK_BINARY_DIR}/CMake/itkTestStreamLongLong.cxx @ONLY IMMEDIATE)
-  IF("ITK_OSTREAM_SUPPORTS_LONG_LONG" MATCHES "^ITK_OSTREAM_SUPPORTS_LONG_LONG$")
-    MESSAGE(STATUS "Checking if ostream supports long long")
-    TRY_COMPILE(ITK_OSTREAM_SUPPORTS_LONG_LONG
-      ${ITK_BINARY_DIR}
-      ${ITK_BINARY_DIR}/CMake/itkTestStreamLongLong.cxx
-      COMPILE_DEFINITIONS -DITK_TEST_OSTREAM_LONG_LONG
-      OUTPUT_VARIABLE OUTPUT)
-    IF(ITK_OSTREAM_SUPPORTS_LONG_LONG)
-      MESSAGE(STATUS "Checking if ostream supports long long -- yes")
-      SET(ITK_OSTREAM_SUPPORTS_LONG_LONG 1 CACHE INTERNAL "Whether ostream supports long long")
-    ELSE(ITK_OSTREAM_SUPPORTS_LONG_LONG)
-      MESSAGE(STATUS "Checking if ostream supports long long -- no")
-      SET(ITK_OSTREAM_SUPPORTS_LONG_LONG 0 CACHE INTERNAL "Whether ostream supports long long")
-      WRITE_FILE(${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeError.log
-        "Determining if ostream supports long long "
-        "failed with the following output:\n"
-        "${OUTPUT}\n" APPEND)
-    ENDIF(ITK_OSTREAM_SUPPORTS_LONG_LONG)
-  ENDIF("ITK_OSTREAM_SUPPORTS_LONG_LONG" MATCHES "^ITK_OSTREAM_SUPPORTS_LONG_LONG$")
-  IF("ITK_ISTREAM_SUPPORTS_LONG_LONG" MATCHES "^ITK_ISTREAM_SUPPORTS_LONG_LONG$")
-    MESSAGE(STATUS "Checking if istream supports long long")
-    TRY_COMPILE(ITK_ISTREAM_SUPPORTS_LONG_LONG
-      ${ITK_BINARY_DIR}
-      ${ITK_BINARY_DIR}/CMake/itkTestStreamLongLong.cxx
-      COMPILE_DEFINITIONS -DITK_TEST_ISTREAM_LONG_LONG
-      OUTPUT_VARIABLE OUTPUT)
-    IF(ITK_ISTREAM_SUPPORTS_LONG_LONG)
-      MESSAGE(STATUS "Checking if istream supports long long -- yes")
-      SET(ITK_ISTREAM_SUPPORTS_LONG_LONG 1 CACHE INTERNAL "Whether istream supports long long")
-    ELSE(ITK_ISTREAM_SUPPORTS_LONG_LONG)
-      MESSAGE(STATUS "Checking if istream supports long long -- no")
-      SET(ITK_ISTREAM_SUPPORTS_LONG_LONG 0 CACHE INTERNAL "Whether istream supports long long")
-      WRITE_FILE(${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeError.log
-        "Determining if istream supports long long "
-        "failed with the following output:\n"
-        "${OUTPUT}\n" APPEND)
-    ENDIF(ITK_ISTREAM_SUPPORTS_LONG_LONG)
-  ENDIF("ITK_ISTREAM_SUPPORTS_LONG_LONG" MATCHES "^ITK_ISTREAM_SUPPORTS_LONG_LONG$")
-ENDIF(ITK_SIZEOF_LONG_LONG)
+#-----------------------------------------------------------------------------
+#
+# Check types such as long long and __int64
+INCLUDE(${ITK_SOURCE_DIR}/CMake/itkTestTypes.cmake)
+
+
+#-----------------------------------------------------------------------------
+# check if the platform has the header file "fenv.h" which has been added as part of the C99 standard
+INCLUDE(CheckIncludeFiles)
+CHECK_INCLUDE_FILES(fenv.h HAVE_FENV_H)
+
+#-----------------------------------------------------------------------------
+# check if the platform has some widely used header files
+CHECK_INCLUDE_FILES("unistd.h" HAVE_UNISTD_H)
 
 #
 # See if compiler preprocessor has the __FUNCTION__ directive used by itkExceptionMacro
@@ -740,6 +696,39 @@
 #
 INCLUDE(${ITK_SOURCE_DIR}/CMake/itkTestFriendTemplatedFunction.cmake)
 
+
+#-----------------------------------------------------------------------------
+# check if the platform has c style IO functions that uses wchar_t * filenames
+# For some reason I can't manage to get a usable result from CHECK_FUNCTION_EXISTS
+TRY_COMPILE(ITK_SUPPORTS_WCHAR_T_FILENAME_CSTYLEIO
+      ${ITK_BINARY_DIR}/CMakeTmp
+      ${ITK_SOURCE_DIR}/CMake/itkTestCStyleIOWCharFilename.cxx
+      OUTPUT_VARIABLE ITK_SUPPORTS_WCHAR_T_FILENAME_CSTYLEIO_OUTPUT)
+
+#-----------------------------------------------------------------------------
+# check if the platform has IO streams constructors that uses wchar_t * filenames
+TRY_COMPILE(ITK_SUPPORTS_WCHAR_T_FILENAME_IOSTREAMS_CONSTRUCTORS
+      ${ITK_BINARY_DIR}/CMakeTmp
+      ${ITK_SOURCE_DIR}/CMake/itkTestIOStreamsWCharFilenameConstructors.cxx
+      OUTPUT_VARIABLE ITK_SUPPORTS_WCHAR_T_FILENAME_IOSTREAMS_CONSTRUCTORS_OUTPUT)
+
+#-----------------------------------------------------------------------------
+# check if the platform can compile fdstream.hpp
+IF(HAVE_UNISTD_H)
+  TRY_COMPILE(ITK_SUPPORTS_FDSTREAM_HPP
+    ${ITK_BINARY_DIR}/CMakeTmp
+    ${ITK_SOURCE_DIR}/CMake/itkTestFDStream.cxx
+    CMAKE_FLAGS -DINCLUDE_DIRECTORIES:STRING=${ITK_SOURCE_DIR}/Utilities/itkExtHdrs
+    COMPILE_DEFINITIONS -DITK_HAVE_UNISTD_H=${HAVE_UNISTD_H}
+    OUTPUT_VARIABLE ITK_SUPPORTS_FDSTREAM_HPP_OUTPUT)
+ELSE(HAVE_UNISTD_H)
+  TRY_COMPILE(ITK_SUPPORTS_FDSTREAM_HPP
+    ${ITK_BINARY_DIR}/CMakeTmp
+    ${ITK_SOURCE_DIR}/CMake/itkTestFDStream.cxx
+    CMAKE_FLAGS -DINCLUDE_DIRECTORIES:STRING=${ITK_SOURCE_DIR}/Utilities/itkExtHdrs
+    OUTPUT_VARIABLE ITK_SUPPORTS_FDSTREAM_HPP_OUTPUT)
+ENDIF(HAVE_UNISTD_H)
+
 #-----------------------------------------------------------------------------
 # Configure the default ITK_DATA_ROOT for the location of ITK Data.
 FIND_PATH(ITK_DATA_ROOT ITKData.readme ${ITK_SOURCE_DIR}/Testing/Data $ENV{ITK_DATA_ROOT})
@@ -915,6 +904,9 @@
 
 find_package(CUDA QUIET REQUIRED)
 
+#set( CUDA_SDK_INCLUDE_PATH /home/ppginf/luis/cuda/common/inc )
+#include_directories( ${CUDA_SDK_INCLUDE_PATH} )
+
 if(NOT CUDA_FOUND)
    message("CUDA could not be located on your system.")
 endif(NOT CUDA_FOUND)
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CMakeLists.txt CITK_CudaCanny//Code/BasicFilters/CMakeLists.txt
--- CITK_lastest_devel_build/Code/BasicFilters/CMakeLists.txt	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CMakeLists.txt	2011-08-07 19:04:20.085186000 -0300
@@ -0,0 +1,37 @@
+SET(CUDA_NVCC_FLAGS 
+  ${CUDA_NVCC_FLAGS}
+  -use_fast_math
+)
+ADD_LIBRARY(ITKBasicFilters itkVTKImageExportBase.cxx)
+CUDA_ADD_LIBRARY(ITKBasicFilters
+CudaCannyEdgeDetection.cu
+CudaDiscreteGaussian.cu
+CudaZeroCrossing.cu
+Cuda2DSeparableConvolution.cu)
+TARGET_LINK_LIBRARIES(ITKBasicFilters ITKCommon)
+IF(ITK_LIBRARY_PROPERTIES)
+  SET_TARGET_PROPERTIES(ITKBasicFilters PROPERTIES ${ITK_LIBRARY_PROPERTIES})
+ENDIF(ITK_LIBRARY_PROPERTIES)
+
+IF(NOT ITK_INSTALL_NO_LIBRARIES)
+  INSTALL(TARGETS ITKBasicFilters
+    RUNTIME DESTINATION ${ITK_INSTALL_BIN_DIR_CM24} COMPONENT RuntimeLibraries
+    LIBRARY DESTINATION ${ITK_INSTALL_LIB_DIR_CM24} COMPONENT RuntimeLibraries
+    ARCHIVE DESTINATION ${ITK_INSTALL_LIB_DIR_CM24} COMPONENT Development)
+ENDIF(NOT ITK_INSTALL_NO_LIBRARIES)
+
+IF(ITK_USE_FLAT_DIRECTORY_INSTALL)
+  SET( ITK_INSTALL_INCLUDE_DIR_CM24_BASICFILTERS ${ITK_INSTALL_INCLUDE_DIR_CM24})
+ELSE(ITK_USE_FLAT_DIRECTORY_INSTALL)
+  SET( ITK_INSTALL_INCLUDE_DIR_CM24_BASICFILTERS ${ITK_INSTALL_INCLUDE_DIR_CM24}/BasicFilters)
+ENDIF(ITK_USE_FLAT_DIRECTORY_INSTALL)
+
+
+IF(NOT ITK_INSTALL_NO_DEVELOPMENT)
+  FILE(GLOB __files1 "${CMAKE_CURRENT_SOURCE_DIR}/*.h")
+  FILE(GLOB __files2 "${CMAKE_CURRENT_SOURCE_DIR}/*.txx")
+  FILE(GLOB __files3 "${CMAKE_CURRENT_SOURCE_DIR}/*.cu")
+  INSTALL(FILES ${__files1} ${__files2} ${__files3}
+    DESTINATION ${ITK_INSTALL_INCLUDE_DIR_CM24_BASICFILTERS}
+    COMPONENT Development)
+ENDIF(NOT ITK_INSTALL_NO_DEVELOPMENT)
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/Cuda2DSeparableConvolution.cu CITK_CudaCanny//Code/BasicFilters/Cuda2DSeparableConvolution.cu
--- CITK_lastest_devel_build/Code/BasicFilters/Cuda2DSeparableConvolution.cu	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/Cuda2DSeparableConvolution.cu	2011-08-07 19:14:40.733186000 -0300
@@ -0,0 +1,147 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: Cuda2DSeparableConvolution.cu,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-09-15 20:58:48 $
+  Version:   $Revision: 4.0.0 $
+
+  Copyright ( c ) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class Cuda2DSeparableConvolutionImageFilter
+ * \brief Class to convolve 2D images
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#include <cuda.h>
+#include <cutil.h>
+
+#include "Cuda2DSeparableConvolution.h"
+
+
+/// allocate texture variables
+texture<float, 1, cudaReadModeElementType> texRef;
+texture<float, 1, cudaReadModeElementType> convTexRef;
+
+
+__global__ void kernel_1DConvolutionH_texture(float *output, int3 size, short halfkernelsize){
+///this version uses the texture memory to store the convolution kernel and the
+///image data
+
+  float sum = 0;
+  int2 pos;
+
+  extern __shared__ float s_mask[];
+
+  ///pixel index of this thread
+  int pixIdx = blockIdx.x * blockDim.x + threadIdx.x;
+
+  ///output pixel index
+  pos.y = __fdividef(pixIdx,size.x);
+  pos.x = pixIdx-(pos.y*size.x);
+
+  if(threadIdx.x<((halfkernelsize<<1)+1)) s_mask[threadIdx.x] = tex1Dfetch(convTexRef,threadIdx.x);
+  
+  __syncthreads();
+
+  for(int k=-halfkernelsize;k<(halfkernelsize+1);k++){
+    sum += (tex1Dfetch(texRef, pixIdx + k * (((pos.x+k)>=0)*((pos.x+k)<size.x))) * s_mask[k+halfkernelsize]);
+  }
+
+  output[pixIdx] = sum;
+}
+
+__global__ void kernel_1DConvolutionV_texture(float *output, int3 size, short halfkernelsize){
+///this version uses the texture memory to store the convolution kernel and the
+///image data
+
+  float sum = 0;
+  int2 pos;
+
+  extern __shared__ float s_mask[];
+
+  ///pixel index of this thread
+  int pixIdx = blockIdx.x * blockDim.x + threadIdx.x;
+
+  ///output pixel index
+  pos.y = __fdividef(pixIdx,size.x);
+  pos.x = pixIdx-(pos.y*size.x);
+
+  if(threadIdx.x<((halfkernelsize<<1)+1)) s_mask[threadIdx.x] = tex1Dfetch(convTexRef,threadIdx.x);
+  
+  __syncthreads();
+
+  for(int k=-halfkernelsize;k<(halfkernelsize+1);k++){
+    sum += (tex1Dfetch(texRef, pixIdx + (size.x*k) * (((pos.y+k)>=0)*((pos.y+k<size.y)))) * s_mask[k+halfkernelsize]);
+  }
+
+  output[pixIdx] = sum;
+}
+
+float* cuda2DSeparableConvolution(dim3 DimGrid, dim3 DimBlock, const float *d_img, int width, int height, const float *d_kernelH, int sizeH, const float *d_kernelV, int sizeV){
+
+  int3 size;
+  size.x = width;
+  size.y = height;
+  size.z = width*height;
+
+  cudaFuncSetCacheConfig(kernel_1DConvolutionH_texture,cudaFuncCachePreferNone);
+  cudaFuncSetCacheConfig(kernel_1DConvolutionV_texture,cudaFuncCachePreferNone);
+
+  short halfKernelWidthH = sizeH >> 1;
+  int kernelSizeH = sizeH*sizeof(float);
+  short halfKernelWidthV = sizeV >> 1;
+  int kernelSizeV = sizeV*sizeof(float);
+
+  float *d_output;
+  cudaMalloc((void**) &d_output, size.z*sizeof(float));
+
+  /// Allocate temporary memory to image data
+  float *d_tmpbuffer;
+  cudaMalloc((void**) &d_tmpbuffer, size.z*sizeof(float));
+
+  /// Bind a texture to the CUDA array
+  cudaBindTexture (NULL, convTexRef, d_kernelH);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// Host side settable texture attributes
+  convTexRef.normalized = false;
+  convTexRef.filterMode = cudaFilterModePoint;
+
+  /// Bind a texture to the CUDA array
+  cudaBindTexture (NULL, texRef, d_img);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// Host side settable texture attributes
+  texRef.normalized = false;
+  texRef.filterMode = cudaFilterModePoint;
+
+  kernel_1DConvolutionH_texture<<<DimGrid,DimBlock,kernelSizeH>>>(d_tmpbuffer,size,halfKernelWidthH);
+
+  /// Bind temporary data texture
+  cudaUnbindTexture(texRef);
+  cudaUnbindTexture(convTexRef);
+  cudaBindTexture (NULL ,texRef, d_tmpbuffer);
+  cudaBindTexture (NULL, convTexRef, d_kernelV);
+
+  kernel_1DConvolutionV_texture<<<DimGrid,DimBlock,kernelSizeV>>>(d_output,size,halfKernelWidthV);
+
+  /// Free allocated memory
+  cudaFree(d_tmpbuffer);
+  cudaUnbindTexture(texRef);
+  cudaUnbindTexture(convTexRef);
+  CUT_CHECK_ERROR("Memory image free failed");
+
+  return(d_output);
+}
+
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/Cuda2DSeparableConvolution.h CITK_CudaCanny//Code/BasicFilters/Cuda2DSeparableConvolution.h
--- CITK_lastest_devel_build/Code/BasicFilters/Cuda2DSeparableConvolution.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/Cuda2DSeparableConvolution.h	2011-08-07 19:14:48.325186000 -0300
@@ -0,0 +1,24 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: Cuda2DSeparableConvolution.h,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-09-20 20:58:48 $
+  Version:   $Revision: 4.0.0 $
+
+  Copyright ( c ) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class Cuda2DSeparableConvolutionImageFilter
+ * \brief Class to convolve 2D images
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+extern "C"
+float* cuda2DSeparableConvolution(dim3 DimGrid, dim3 DimBlock, const float *d_img, int width, int height, const float *d_kernelH, int sizeH, const float *d_kernelV, int sizeV);
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaCannyEdgeDetection.cu CITK_CudaCanny//Code/BasicFilters/CudaCannyEdgeDetection.cu
--- CITK_lastest_devel_build/Code/BasicFilters/CudaCannyEdgeDetection.cu	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaCannyEdgeDetection.cu	2011-08-07 19:15:06.285186001 -0300
@@ -0,0 +1,432 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: CudaCannyEdgeDetection.cu,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-05-20 12:01:33 $
+  Version:   $Revision: 2.0.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class CannyEdgeDetectionImageFilter
+ * \brief CUDA implementation of Canny's algorithm
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#include <cuda.h>
+#include <cutil.h>
+
+#include "CudaCannyEdgeDetection.h"
+
+
+/// allocate texture variables
+texture<float, 1, cudaReadModeElementType> texRef;
+texture<float, 1, cudaReadModeElementType> mag_texRef;
+texture<float, 1, cudaReadModeElementType> hysTexRef;
+texture<float, 1, cudaReadModeElementType> der_texRef;
+
+
+__global__ void kernel_Compute2ndDerivativePos(float *Magnitude, int3 size){
+/// This kernel receive the blurred image and it's second derivative and returns
+/// the Gradient Magnitude of the image.
+/// It ignores completely the borders of the image because the border's pixels
+/// doesn't have all neighbors.
+
+  /// Pixel index of this thread
+  int pixIdx = blockIdx.x * blockDim.x + threadIdx.x;
+
+  /// Output pixel index
+  int2 pos;
+  pos.y = __fdividef(pixIdx,size.x);
+  pos.x = pixIdx-(pos.y*size.x);
+
+  float4 cross_Lvv;
+  float4 cross_L;
+
+  float Lx = 0,
+        Ly = 0,
+        Lvvx = 0,
+        Lvvy = 0,
+        gradMag;
+
+  cross_L.x = tex1Dfetch(texRef,(pixIdx-(size.x*(pos.y>0))));
+  cross_L.y = tex1Dfetch(texRef,(pixIdx+(size.x*(pos.y<(size.y-1)))));
+  cross_L.z = tex1Dfetch(texRef,(pixIdx-(pos.x>0)));
+  cross_L.w = tex1Dfetch(texRef,(pixIdx+(pos.x<(size.x-1))));
+  cross_Lvv.x = tex1Dfetch(der_texRef,(pixIdx-(size.x*(pos.y>0))));
+  cross_Lvv.y = tex1Dfetch(der_texRef,(pixIdx+(size.x*(pos.y<(size.y-1)))));
+  cross_Lvv.z = tex1Dfetch(der_texRef,(pixIdx-(pos.x>0)));
+  cross_Lvv.w = tex1Dfetch(der_texRef,(pixIdx+(pos.x<(size.x-1))));
+    
+  Lx = (-0.5*cross_L.z) + (0.5*cross_L.w);
+  Ly = (0.5*cross_L.x) - (0.5*cross_L.y);
+
+  Lvvx = (-0.5*cross_Lvv.z) + (0.5*cross_Lvv.w);
+  Lvvy = (0.5*cross_Lvv.x) - (0.5*cross_Lvv.y);
+
+  gradMag = sqrt((Lx*Lx)+(Ly*Ly));
+
+  Magnitude[pixIdx] = (((Lvvx*(Lx/gradMag)+Lvvy*(Ly/gradMag))<=0)*gradMag);
+
+}
+
+float* cuda2ndDerivativePos(dim3 DimGrid, dim3 DimBlock, const float *d_input, const float *d_Lvv, int width, int height){
+
+  int3 size;
+  size.x = width;
+  size.y = height;
+  size.z = width*height;
+
+  cudaFuncSetCacheConfig(kernel_Compute2ndDerivativePos,cudaFuncCachePreferNone);
+
+ /// Allocate output memory to image data
+  float * d_mag;
+  cudaMalloc((void**) &d_mag, (size.z*sizeof(float)));
+  CUT_CHECK_ERROR("Memory image creation failed");
+
+  /// bind a texture to the CUDA array
+  cudaBindTexture (NULL, texRef, d_input);
+  cudaBindTexture (NULL, der_texRef, d_Lvv);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// host side settable texture attributes
+  texRef.normalized = false;
+  texRef.filterMode = cudaFilterModePoint;
+
+  kernel_Compute2ndDerivativePos<<<DimGrid,DimBlock>>>(d_mag, size);
+
+  /// unbind texture reference
+  cudaUnbindTexture(texRef);
+  cudaUnbindTexture(der_texRef);
+  CUT_CHECK_ERROR("Memory image free failed");
+
+  return(d_mag);
+
+}
+
+
+__global__ void kernel_Compute2ndDerivative(float *Lvv, int3 size){
+/// This kernel receives a blurred image and returns it's second derivative.
+/// It ignores completely the borders of the image becausa the border's pixels
+/// doesn't have all neighbors
+
+  /// Pixel index of this thread
+  int pixIdx = blockIdx.x * blockDim.x + threadIdx.x;
+
+  /// Output pixel index
+  int2 pos;
+  pos.y = __fdividef(pixIdx,size.x);
+  pos.x = pixIdx-(pos.y*size.x);
+
+  float4 diagonal;
+  float4 cross;
+  float pixel = tex1Dfetch(texRef,(pixIdx));
+
+  float Lx = 0, 
+        Ly = 0, 
+        Lxx = 0, 
+        Lxy = 0, 
+        Lyy = 0;
+
+  /// Stores the neighbors of the pixel on variables, because they will be
+  /// readen more than one time.
+  diagonal.x = tex1Dfetch(texRef,(pixIdx+((-size.x-1)*(pos.y>0)*(pos.x>0))));
+  diagonal.y = tex1Dfetch(texRef,(pixIdx+((-size.x+1)*(pos.y>0)*(pos.x<(size.x-1)))));
+  diagonal.z = tex1Dfetch(texRef,(pixIdx+((size.x-1)*(pos.y<(size.y-1))*(pos.x>0))));
+  diagonal.w = tex1Dfetch(texRef,(pixIdx+((size.x+1)*(pos.y<(size.y-1))*(pos.x<(size.x-1)))));
+  cross.x = tex1Dfetch(texRef,(pixIdx-(size.x*(pos.y>0))));
+  cross.y = tex1Dfetch(texRef,(pixIdx+(size.x*(pos.y<(size.y-1)))));
+  cross.z = tex1Dfetch(texRef,(pixIdx-(pos.x>0)));
+  cross.w = tex1Dfetch(texRef,(pixIdx+(pos.x<(size.x-1))));
+
+  Lx = (-0.5*cross.z) + (0.5*cross.w);
+  Ly = (0.5*cross.x) - (0.5*cross.y);
+  Lxx = cross.z - (2.0*pixel) + cross.w;
+  Lxy = (-0.25*diagonal.x) + (0.25*diagonal.y) + (0.25*diagonal.z) + (-0.25*diagonal.w);
+  Lyy = cross.x -(2.0*pixel) + cross.y;
+
+
+  Lvv[pixIdx] = (((Lx*Lx)*Lxx) + (2.0*Lx*Ly*Lxy) + (Ly*Ly*Lyy))/((Lx*Lx) + (Ly*Ly));
+
+}
+
+float* cuda2ndDerivative(dim3 DimGrid, dim3 DimBlock, const float *d_input, int width, int height){
+
+  int3 size;
+  size.x = width;
+  size.y = height;
+  size.z = width*height;
+
+  cudaFuncSetCacheConfig(kernel_Compute2ndDerivative,cudaFuncCachePreferNone);
+
+ /// Allocate output memory to image data
+  float * d_Lvv;
+  cudaMalloc((void**) &d_Lvv, (size.z*sizeof(float)));
+  CUT_CHECK_ERROR("Memory image creation failed");
+
+  /// bind a texture to the CUDA array
+  cudaBindTexture (NULL ,texRef, d_input);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// host side settable texture attributes
+  texRef.normalized = false;
+  texRef.filterMode = cudaFilterModePoint;
+
+  kernel_Compute2ndDerivative<<<DimGrid,DimBlock>>>(d_Lvv, size);
+
+  /// unbind texture reference
+  cudaUnbindTexture(texRef);
+  CUT_CHECK_ERROR("Memory image free failed");
+
+  return(d_Lvv);
+
+}
+
+
+__global__ void hysteresisPreparation_kernel(float *hysteresis, int3 size, float t1, float t2){
+
+  ///pixel index of this thread
+  int pixIdx = blockIdx.x * blockDim.x + threadIdx.x;
+  float pixel;
+
+
+  ///output pixel index
+  int2 pos;
+  pos.y = __fdividef(pixIdx,size.x);
+  pos.x = pixIdx-(pos.y*size.x);
+
+  if(pixIdx < size.z){
+
+    pixel = tex1Dfetch(mag_texRef, pixIdx)*tex1Dfetch(texRef,pixIdx);
+
+    hysteresis[pixIdx] = ((POSSIBLE_EDGE-1)*(pixel>t2)+POSSIBLE_EDGE)*(pixel>t1);
+
+  }
+
+}
+
+__global__ void hysteresisWrite_kernel(float *output, int3 size){
+
+  ///pixel index of this thread
+  int pixIdx = blockIdx.x * blockDim.x + threadIdx.x;
+  float pixel;
+
+  if(pixIdx < size.z){
+
+    pixel = tex1Dfetch(hysTexRef, pixIdx);
+    output[pixIdx] = (pixel==DEFINITIVE_EDGE) * pixel;
+
+  }
+
+}
+
+__global__ void kernel_hysteresis_glm1D(float *hys_img, int3 size, int *have_modified_pixels){
+
+  __shared__ float s_slice[SLICE_WIDTH*SLICE_WIDTH];
+  __shared__ int modified_block_pixels; /// control inner loop
+  __shared__ int modified_image_pixels; /// control outer loop
+  int edge;
+
+  // pixel position indexes on block
+  int2 slice_pos;
+  slice_pos.y = threadIdx.x / SLICE_BLOCK_WIDTH;
+  slice_pos.x = threadIdx.x - (slice_pos.y * SLICE_BLOCK_WIDTH);
+
+  // pixel positions indexes on image
+  int2 pos;
+  pos.x = (slice_pos.x + (blockIdx.x * SLICE_BLOCK_WIDTH)) % size.x;
+  pos.y = (((((blockIdx.x * SLICE_BLOCK_WIDTH))) / size.x) * SLICE_BLOCK_WIDTH ) + slice_pos.y;
+
+  int sliceIdx = (slice_pos.y*SLICE_WIDTH) + slice_pos.x  + SLICE_WIDTH + 1;
+
+  // pixel position at the hysteresis image
+  int pixIdx = pos.y * size.x + pos.x;
+
+  if (!threadIdx.x) modified_image_pixels = NOT_MODIFIED;
+  if (!pixIdx) have_modified_pixels[0] = 0;
+  
+
+  // load center
+  if ((pos.x>0)&&(pos.y>0)&&(pos.x<(size.x-1))&&(pos.y<(size.y-1))){
+
+    s_slice[sliceIdx] = hys_img[pixIdx];
+
+    /// load top
+    if(!slice_pos.y){
+      if(!slice_pos.x){
+        s_slice[0] = hys_img[pixIdx-size.x-1];///<TL
+      }
+      s_slice[slice_pos.x+1] = hys_img[pixIdx-size.x];
+      if(slice_pos.x == (SLICE_BLOCK_WIDTH-1)){
+        s_slice[SLICE_WIDTH-1] = hys_img[pixIdx-size.x+1];///<TR
+      }
+    }
+    /// load bottom
+    if(slice_pos.y == (SLICE_BLOCK_WIDTH-1)){
+      if(!slice_pos.x){
+        s_slice[SLICE_WIDTH*(SLICE_WIDTH-1)] = hys_img[pixIdx+size.x-1];///<BL
+      }
+      s_slice[(SLICE_WIDTH*(SLICE_WIDTH-1))+1+slice_pos.x] = hys_img[pixIdx+size.x];
+      if(slice_pos.x == (SLICE_BLOCK_WIDTH-1)){
+        s_slice[(SLICE_WIDTH*SLICE_WIDTH)-1] = hys_img[pixIdx+size.x+1];///<BR
+      }
+    }
+    /// load left
+    if(!slice_pos.x){
+      s_slice[(slice_pos.y+1)*SLICE_WIDTH] = hys_img[pixIdx-1];
+    }
+    /// load right
+    if(slice_pos.x == (SLICE_BLOCK_WIDTH-1)){
+      s_slice[((slice_pos.y+2)*SLICE_WIDTH)-1] = hys_img[pixIdx+1];
+    }
+
+  }
+  else{
+    s_slice[sliceIdx] = 0;
+
+    /// load top
+    if(!slice_pos.y){
+      if(!slice_pos.x){
+        s_slice[0] = 0;///<TL
+      }
+      s_slice[slice_pos.x+1] = 0;
+      if(slice_pos.x == (SLICE_BLOCK_WIDTH-1)){
+        s_slice[SLICE_WIDTH-1] = 0;///<TR
+      }
+    }
+    /// load bottom
+    if(slice_pos.y == (SLICE_BLOCK_WIDTH-1)){
+      if(!slice_pos.x){
+        s_slice[SLICE_WIDTH*(SLICE_WIDTH-1)] = 0;///<BL
+      }
+      s_slice[(SLICE_WIDTH*(SLICE_WIDTH-1))+1+slice_pos.x] = 0;
+      if(slice_pos.x == (SLICE_BLOCK_WIDTH-1)){
+        s_slice[(SLICE_WIDTH*SLICE_WIDTH)-1] = 0;///<BR
+      }
+    }
+    /// load left
+    if(!slice_pos.x){
+      s_slice[(slice_pos.y+1)*SLICE_WIDTH] = 0;
+    }
+    /// load right
+    if(slice_pos.x == (SLICE_BLOCK_WIDTH-1)){
+      s_slice[((slice_pos.y+2)*SLICE_WIDTH)-1] = 0;
+    }
+  }
+
+  __syncthreads();
+
+  do{
+
+    if (!threadIdx.x) modified_block_pixels = NOT_MODIFIED;
+
+    __syncthreads();
+
+    if(s_slice[sliceIdx] == POSSIBLE_EDGE){
+
+      /// edge = 1 if at last one pixel's neighbour is a definitive edge 
+      /// and edge = 0 if doesn't
+      edge = (!((s_slice[sliceIdx-SLICE_WIDTH-1] != DEFINITIVE_EDGE) *
+                (s_slice[sliceIdx-SLICE_WIDTH] != DEFINITIVE_EDGE) *
+                (s_slice[sliceIdx-SLICE_WIDTH+1] != DEFINITIVE_EDGE) *
+                (s_slice[sliceIdx-1] != DEFINITIVE_EDGE) *
+                (s_slice[sliceIdx+1] != DEFINITIVE_EDGE) *
+                (s_slice[sliceIdx+SLICE_WIDTH-1] != DEFINITIVE_EDGE) *
+                (s_slice[sliceIdx+SLICE_WIDTH] != DEFINITIVE_EDGE) *
+                (s_slice[sliceIdx+SLICE_WIDTH+1] != DEFINITIVE_EDGE)));
+      if ( (edge)*(!modified_block_pixels) ) modified_image_pixels = modified_block_pixels = MODIFIED;
+      s_slice[sliceIdx] = POSSIBLE_EDGE + ((edge)*(POSSIBLE_EDGE-1));
+
+    }
+
+    __syncthreads();
+
+
+  }while(modified_block_pixels);// end inner loop
+
+  // Update only the blocks that have modified pixels
+  if ((modified_image_pixels) && (pos.x < (size.x)) && (pos.y < (size.y))){
+    hys_img[pixIdx] = s_slice[sliceIdx];
+  }
+  if (!threadIdx.x) have_modified_pixels[0] += modified_image_pixels;
+
+}
+
+float* cudaHysteresis(dim3 DimGrid, dim3 DimBlock, float *d_img, float *d_mag, int width, int height, float t1, float t2){
+ 
+  int3 size;
+  size.x = width;
+  size.y = height;
+  size.z = width*height;
+
+  cudaFuncSetCacheConfig(hysteresisPreparation_kernel,cudaFuncCachePreferNone);
+  cudaFuncSetCacheConfig(kernel_hysteresis_glm1D,cudaFuncCachePreferNone);
+  cudaFuncSetCacheConfig(hysteresisWrite_kernel,cudaFuncCachePreferNone);
+
+  float *d_hys;
+  cudaMalloc((void**) &d_hys, (size.z*sizeof(float)));
+  CUT_CHECK_ERROR("Memory hysteresis image creation failed");
+
+  /// bind a texture to the image
+  cudaBindTexture (NULL ,texRef, d_img);
+  cudaBindTexture (NULL ,mag_texRef, d_mag);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// host side settable texture attributes
+  texRef.normalized = false;
+  texRef.filterMode = cudaFilterModePoint;
+ 
+  hysteresisPreparation_kernel<<<DimGrid,DimBlock>>>(d_hys, size, t1, t2);
+
+  /// free allocated memory
+  cudaUnbindTexture(texRef);
+  cudaUnbindTexture(mag_texRef);
+  CUT_CHECK_ERROR("Memory unbind failed");
+
+  // counter of modifications
+  int *modif;
+  cudaMalloc((void**) &modif, (sizeof(int)));
+
+  int cont[1];
+  do{
+    kernel_hysteresis_glm1D<<<DimGrid,DimBlock>>>(d_hys, size, modif);
+    CUT_CHECK_ERROR("Hysteresis Kernel failed");
+    cudaMemcpy(cont,modif,sizeof(int),cudaMemcpyDeviceToHost);
+  }while(cont[0]);
+
+  /// bind a texture to the CUDA array
+  cudaBindTexture (NULL ,hysTexRef, d_hys);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// host side settable texture attributes
+  hysTexRef.normalized = false;
+  hysTexRef.filterMode = cudaFilterModePoint;
+
+  hysteresisWrite_kernel<<<DimGrid,DimBlock>>>(d_img, size);
+  CUT_CHECK_ERROR("Hysteresis Write failed");
+
+  /// free allocated memory
+  cudaUnbindTexture(hysTexRef);
+  CUT_CHECK_ERROR("Memory unbind failed");
+
+  cudaFree(d_hys);
+  cudaFree(modif);
+  CUT_CHECK_ERROR("Memory free failed");
+
+  //returna edges  
+  float *d_edges;
+  d_edges = d_img;
+  return(d_edges);
+
+}
+
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaCannyEdgeDetection.h CITK_CudaCanny//Code/BasicFilters/CudaCannyEdgeDetection.h
--- CITK_lastest_devel_build/Code/BasicFilters/CudaCannyEdgeDetection.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaCannyEdgeDetection.h	2011-08-07 19:15:06.349186000 -0300
@@ -0,0 +1,43 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: CudaCannyEdgeDetection.h,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-05-10 12:01:33 $
+  Version:   $Revision: 2.2.1 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class CannyEdgeDetectionImageFilter
+ * \brief CUDA implementation of Canny's algorithm
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+// pixel types used on hysteresis
+#define DEFINITIVE_EDGE 255
+#define POSSIBLE_EDGE 128
+#define NON_EDGE 0
+
+// sizes for hysteresis' slice width
+#define SLICE_WIDTH 18
+#define SLICE_BLOCK_WIDTH 16
+
+// there's a pixel modified or not on the slice
+#define MODIFIED 1
+#define NOT_MODIFIED 0
+
+extern "C"
+float* cuda2ndDerivativePos(dim3 DimGrid, dim3 DimBlock, const float *d_input, const float *d_Lvv, int width, int height);
+
+extern "C"
+float* cuda2ndDerivative(dim3 DimGrid, dim3 DimBlock, const float *d_input, int width, int height);
+
+extern "C"
+float* cudaHysteresis(dim3 DimGrid, dim3 DimBlock, float *d_img, float *d_gauss, int width, int height, float t1, float t2);
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaDiscreteGaussian.cu CITK_CudaCanny//Code/BasicFilters/CudaDiscreteGaussian.cu
--- CITK_lastest_devel_build/Code/BasicFilters/CudaDiscreteGaussian.cu	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaDiscreteGaussian.cu	2011-08-07 19:15:24.701186001 -0300
@@ -0,0 +1,82 @@
+/*=========================================================================
+
+Program:   Insight Segmentation & Registration Toolkit
+Module:    $RCSfile: CudaDiscreteGaussian.cu,v $
+Language:  C/C++ CUDA
+Date:      $Date: 2010-09-15 12:44:26 $
+Version:   $Revision: 3.1.0 $
+
+Copyright (c) Insight Software Consortium. All rights reserved.
+See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+This software is distributed WITHOUT ANY WARRANTY; without even 
+the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/**
+ * \class CudaDiscreteGaussianImageFilter
+ * \brief Blurs an image by separable convolution with discrete gaussian kernels.
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#include <cuda.h>
+#include <cutil.h>
+
+#include "CudaDiscreteGaussian.h"
+
+/// allocate texture variables
+texture<float, 1, cudaReadModeElementType> texRef;
+texture<float, 1, cudaReadModeElementType> gaussTexRef;
+
+
+
+__global__ void calculateGaussianKernel(float *gaussKernel, const float sigma, int halfKernelWidth){
+
+  /// pixel index of this thread
+  /// this makes the normal curve
+  int i = threadIdx.x - halfKernelWidth;
+  extern __shared__ float s_gaussKernel[];
+  __shared__ float sum;
+  
+  /// this kernel must allocate 'kernelWidth' threads
+  s_gaussKernel[threadIdx.x] = (__fdividef(1,(sqrtf(2*M_PI*sigma))))*expf((-1)*(__fdividef((i*i),(2*sigma*sigma))));
+
+  __syncthreads();
+
+  /// Thread 0 sum all the gassian kernel array
+  // This is not so bad because the array is always short
+  if (!threadIdx.x) {
+    int th;
+    sum = 0;
+    for(th = 0; th<blockDim.x; th++) sum += s_gaussKernel[th];
+  }
+
+  __syncthreads();
+
+  gaussKernel[threadIdx.x] = s_gaussKernel[threadIdx.x]/sum;
+
+}
+
+float* cuda1DGaussianOperator(dim3 DimGrid, dim3 DimBlock, unsigned int width, float gaussianVariance){
+
+  /// The Gaussian Kernel Width must be odd
+  if (width < 1) width = 1;
+  if (width%2 == 0) width--;
+  short halfWidth = width >> 1;
+
+  int kernelSize = width*sizeof(float);
+
+  float *cudaGaussKernel;
+  cudaMalloc((void**)&cudaGaussKernel,kernelSize);
+
+  /// Calculate gaussian kernel
+  calculateGaussianKernel<<<DimGrid,DimBlock,kernelSize>>>(cudaGaussKernel, gaussianVariance, halfWidth);
+  
+  return(cudaGaussKernel);
+}
+
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaDiscreteGaussian.h CITK_CudaCanny//Code/BasicFilters/CudaDiscreteGaussian.h
--- CITK_lastest_devel_build/Code/BasicFilters/CudaDiscreteGaussian.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaDiscreteGaussian.h	2011-08-07 19:15:24.765186000 -0300
@@ -0,0 +1,25 @@
+/*=========================================================================
+
+Program:   Insight Segmentation & Registration Toolkit
+Module:    $RCSfile: CudaDiscreteGaussian.h,v $
+Language:  C/C++ CUDA
+Date:      $Date: 2010-09-10 12:44:26 $
+Version:   $Revision: 3.2.0 $
+
+Copyright (c) Insight Software Consortium. All rights reserved.
+See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+This software is distributed WITHOUT ANY WARRANTY; without even 
+the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/**
+ * \class CudaDiscreteGaussianImageFilter
+ * \brief Blurs an image by separable convolution with discrete gaussian kernels.
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+extern "C"
+float* cuda1DGaussianOperator(dim3 DimGrid, dim3 DimBlock, unsigned int width, float gaussianVariance);
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaSobelEdgeDetection.cu CITK_CudaCanny//Code/BasicFilters/CudaSobelEdgeDetection.cu
--- CITK_lastest_devel_build/Code/BasicFilters/CudaSobelEdgeDetection.cu	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaSobelEdgeDetection.cu	2011-08-07 19:15:45.781186000 -0300
@@ -0,0 +1,120 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: CudaSobelEdgeDetection.cu,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-09-10 20:50:03 $
+  Version:   $Revision: 3.0.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/**
+ * \class SobelEdgeDetectionImageFilter
+ * \brief A 2D edge detection using the Sobel operator implemented in CUDA.
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#include <cuda.h>
+#include <cutil.h>
+
+#include "CudaSobelEdgeDetection.h"
+
+#define THREADS_PER_BLOCK 256
+
+/// allocate texture variables
+texture<float, 1, cudaReadModeElementType> texRef;
+
+
+__global__ void kernel_2DSobel(float *Magnitude, float *Direction, int3 size){
+/// This is an less elaborated kernel version that calculate the sobel-x, sobel-y, 
+/// then uses the calculated values to return to memory just the needed information.
+/// That is, the magnitude and direction of the edge.
+/// It ignores completely the borders of the image becausa the border's pixels
+/// doesn't have all neighbors
+
+  float2 g_i;
+  g_i.x = g_i.y = 0;
+  float4 diagonal;
+  float4 cross;
+
+  /// Pixel index of this thread
+  int pixIdx = blockIdx.x * blockDim.x + threadIdx.x;
+
+  /// Output pixel index
+  int2 pos;
+  pos.y = __fdividef(pixIdx,size.x);
+  pos.x = pixIdx-(pos.y*size.x);
+
+  ///Ignore the image borders
+  if ((pos.x) && ((size.x-1)-pos.x) && (pos.y) && ((size.y-1)-pos.y)){
+
+    /// Stores the neighbors of the pixel on variables, because they will be
+    /// readen more than one time.
+    diagonal.x = tex1Dfetch(texRef,(pixIdx-size.x-1));
+    diagonal.y = tex1Dfetch(texRef,(pixIdx-size.x+1));
+    diagonal.z = tex1Dfetch(texRef,(pixIdx+size.x-1));
+    diagonal.w = tex1Dfetch(texRef,(pixIdx+size.x+1));
+    cross.x = tex1Dfetch(texRef,(pixIdx-size.x));
+    cross.y = tex1Dfetch(texRef,(pixIdx+size.x));
+    cross.z = tex1Dfetch(texRef,(pixIdx-1));
+    cross.w = tex1Dfetch(texRef,(pixIdx+1));
+
+    /// SobelX
+    g_i.x -= (diagonal.x+cross.z+cross.z+diagonal.z);
+    g_i.x += (diagonal.y+cross.w+cross.w+diagonal.w);
+    
+    /// SobelY
+    g_i.y -= (diagonal.z+cross.y+cross.y+diagonal.w);
+    g_i.y += (diagonal.x+cross.x+cross.x+diagonal.y);
+    
+  }
+
+  Magnitude[pixIdx] = sqrtf((g_i.x*g_i.x) + (g_i.y*g_i.y));
+
+  Direction[pixIdx] = (g_i.x != 0)*(atanf(__fdividef(g_i.y,g_i.x)));
+
+}
+
+extern "C"
+Tgrad* cudaSobel(Tgrad *d_gradient, const float *d_img, int width, int height){
+
+  int3 size;
+  size.x = width;
+  size.y = height;
+  size.z = width*height;
+
+  int blocksPerGrid = ((size.z) + THREADS_PER_BLOCK -1)/THREADS_PER_BLOCK;
+  dim3 DimBlock(THREADS_PER_BLOCK,1,1);
+  dim3 DimGrid(blocksPerGrid,1,1);
+
+  /// Allocate output memory to image data
+  cudaMalloc((void**) &d_gradient->Magnitude, size.z*sizeof(float));
+  cudaMalloc((void**) &d_gradient->Direction, size.z*sizeof(float));
+
+  /// bind a texture to the CUDA array
+  cudaBindTexture (NULL ,texRef, d_img);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// host side settable texture attributes
+  texRef.normalized = false;
+  texRef.filterMode = cudaFilterModePoint;
+
+  kernel_2DSobel<<<DimGrid,DimBlock>>>(d_gradient->Magnitude, d_gradient->Direction, size);
+
+  /// unbind texture reference
+  cudaUnbindTexture(texRef);
+  CUT_CHECK_ERROR("Memory image free failed");
+
+  return(d_gradient);
+}
+
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaSobelEdgeDetection.h CITK_CudaCanny//Code/BasicFilters/CudaSobelEdgeDetection.h
--- CITK_lastest_devel_build/Code/BasicFilters/CudaSobelEdgeDetection.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaSobelEdgeDetection.h	2011-08-07 19:15:45.845186000 -0300
@@ -0,0 +1,37 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: CudaSobelEdgeDetection.h,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-09-22 20:50:03 $
+  Version:   $Revision: 3.0.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/**
+ * \class SobelEdgeDetectionImageFilter
+ * \brief A 2D edge detection using the Sobel operator implemented in CUDA.
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+
+template <class TMagnitude, class TDirection>
+class Gradient
+{
+public:
+  TMagnitude *Magnitude;
+  TDirection *Direction;
+};
+
+typedef Gradient<float,float>  Tgrad;
+
+extern "C"
+Tgrad* cudaSobel(Tgrad* d_gradient, const float *d_img, int width, int height);
+
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaZeroCrossing.cu CITK_CudaCanny//Code/BasicFilters/CudaZeroCrossing.cu
--- CITK_lastest_devel_build/Code/BasicFilters/CudaZeroCrossing.cu	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaZeroCrossing.cu	2011-08-07 19:16:07.309186000 -0300
@@ -0,0 +1,102 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: CudaZeroCrossing.cu,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-10-26 17:31:03 $
+  Version:   $Revision: 0.0.1 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class CudaZeroCrossingImageFilter
+ * \brief This filter finds the closest pixel to the zero-crossings in a Image. 
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#include <cuda.h>
+#include <cutil.h>
+
+#include "CudaZeroCrossing.h"
+
+
+/// allocate texture variables
+texture<float, 1, cudaReadModeElementType> texRef;
+
+__global__ void kernel_zerocrossing(float* image, int3 size){
+///this is the kernel to calculate the zero-crossing of the image. It is
+///implemmented using texture fetching. The dependence of inter-block data makes
+///the use of shared memory hard-boiled.
+
+  float  pixel;
+  float4 cross;
+  float res = 0;
+
+  int pixIdx = blockDim.x * blockIdx.x + threadIdx.x;
+  pixel = tex1Dfetch(texRef,pixIdx);
+
+  ///output pixel index
+  int2 pos;
+  pos.y = __fdividef(pixIdx,size.x);
+  pos.x = pixIdx-(pos.y*size.x);
+
+  cross.x = tex1Dfetch(texRef,pixIdx-(pos.x>0));
+  cross.y = tex1Dfetch(texRef,pixIdx-(size.x*(pos.y>0)));
+  cross.z = tex1Dfetch(texRef,pixIdx+(pos.x<(size.x-1)));
+  cross.w = tex1Dfetch(texRef,pixIdx+(size.x*(pos.y<(size.y-1))));
+
+  res = (((pixel*cross.x)<=0) *\
+      (fabsf(pixel) < fabsf(cross.x)));
+
+  res = res || ((((pixel*cross.y)<=0)) *\
+      (fabsf(pixel) < fabsf(cross.y)));
+    
+  res = res || ((((pixel*cross.z)<=0)) *\
+      (fabsf(pixel) <= fabsf(cross.z)));
+
+  res = res || ((((pixel*cross.w)<=0)) *\
+      (fabsf(pixel) <= fabsf(cross.w)));
+
+  image[pixIdx] = res;
+  
+}
+
+float* cudaZeroCrossing(dim3 DimGrid, dim3 DimBlock, float *d_input, int width, int height){
+
+  int3 size;
+  size.x = width;
+  size.y = height;
+  size.z = width*height;
+
+  cudaFuncSetCacheConfig(kernel_zerocrossing,cudaFuncCachePreferNone);
+
+  float *d_img;
+  cudaMalloc((void**) &d_img, (size.z*sizeof(float)));
+  CUT_CHECK_ERROR("Memory image creation failed");
+
+  /// bind a texture to the CUDA array
+  cudaBindTexture (NULL, texRef, d_input);
+  CUT_CHECK_ERROR("Texture bind failed");
+
+  /// host side settable texture attributes
+  texRef.normalized = false;
+  texRef.filterMode = cudaFilterModePoint;
+  
+  kernel_zerocrossing<<<DimGrid,DimBlock>>>(d_img, size);
+  
+  ///free allocated memory
+  cudaUnbindTexture(texRef);
+  CUT_CHECK_ERROR("Memory image free failed");
+ 
+  return(d_img);
+
+}
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/CudaZeroCrossing.h CITK_CudaCanny//Code/BasicFilters/CudaZeroCrossing.h
--- CITK_lastest_devel_build/Code/BasicFilters/CudaZeroCrossing.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/CudaZeroCrossing.h	2011-08-07 19:16:07.365186000 -0300
@@ -0,0 +1,25 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: CudaZeroCrossing.h,v $
+  Language:  C/C++ CUDA
+  Date:      $Date: 2010-10-26 17:31:03 $
+  Version:   $Revision: 0.0.1 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class CudaZeroCrossingImageFilter
+ * \brief This filter finds the closest pixel to the zero-crossings in a Image. 
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+extern "C"
+float* cudaZeroCrossing(dim3 DimGrid, dim3 DimBlock, float *d_input, int width, int height);
+
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.h CITK_CudaCanny//Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.h
--- CITK_lastest_devel_build/Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.h	2011-08-07 19:16:32.277186001 -0300
@@ -0,0 +1,102 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCuda2DSeparableConvolutionImageFilter.h,v $
+  Language:  C++
+  Date:      $Date: 2010-11-11 20:58:48 $
+  Version:   $Revision: 0.1 $
+
+  Copyright ( c ) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class Cuda2DSeparableConvolutionImageFilter
+ * \brief Class to convolve 2D images
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#ifndef __itkCuda2DSeparableConvolutionImageFilter_h
+#define __itkCuda2DSeparableConvolutionImageFilter_h
+
+#include "itkImageToImageFilter.h"
+#include "itkCudaInterface.h"
+
+#include "cuda.h"
+#include "Cuda2DSeparableConvolution.h"
+
+namespace itk {
+
+template<class TInputImage, class TOutputImage>
+class ITK_EXPORT Cuda2DSeparableConvolutionImageFilter
+: public ImageToImageFilter<TInputImage, TOutputImage>
+{
+public:
+  typedef Cuda2DSeparableConvolutionImageFilter                Self;
+  typedef ImageToImageFilter<TInputImage, TOutputImage>        Superclass;
+  typedef SmartPointer<Self>                                   Pointer;
+  typedef SmartPointer<const Self>                             ConstPointer;
+
+  /** Method for creation through the object factory. */
+  itkNewMacro( Self );
+
+  /** Run-time type information ( and related methods ) */
+  itkTypeMacro( Cuda2DSeparableConvolutionImageFilter, ImageToImageFilter );
+
+  /** Dimensionality of input and output data is assumed to be the same. */
+  itkStaticConstMacro( ImageDimension, unsigned int,
+                       TInputImage::ImageDimension );
+
+
+  typedef TInputImage                                    InputImageType;
+  typedef TOutputImage                                   OutputImageType;
+  typedef typename InputImageType::PixelType             InputPixelType;
+  typedef typename OutputImageType::PixelType            OutputPixelType;
+  typedef typename OutputImageType::RegionType           OutputRegionType;
+
+  /** Cuda2DSeparableConvolutionImageFilter needs two smaller extra inputs 
+   * (the image kernels to be aplyed on each direction)
+   * requested region than output requested region.  As such, this filter
+   * needs to provide an implementation for GenerateInputRequestedRegion() in
+   * order to inform the pipeline execution model.
+   * \sa ProcessObject::GenerateInputRequestedRegion()  */
+//  virtual void GenerateInputRequestedRegion();
+
+  void SetInputMaskHorizontal( OutputPixelType * , unsigned int);
+  const OutputPixelType * GetInputMaskHorizontal( void );
+  void SetInputMaskVertical( OutputPixelType * , unsigned int);
+  const OutputPixelType * GetInputMaskVertical( void );
+
+protected:
+  /** de/constructor */
+  Cuda2DSeparableConvolutionImageFilter();
+  ~Cuda2DSeparableConvolutionImageFilter();
+
+  void GenerateData();
+
+  typedef CudaInterface CudaInterfaceType;
+  
+private:
+  Cuda2DSeparableConvolutionImageFilter( const Self& ); //purposely not implemented
+  void operator=( const Self& ); //purposely not implemented
+
+  typename CudaInterfaceType::Pointer m_CudaConf;
+private:
+
+  OutputPixelType * m_Mask1;
+  unsigned int      m_SizeMask1;
+  OutputPixelType * m_Mask2;
+  unsigned int      m_SizeMask2;
+
+};
+}
+
+#ifndef ITK_MANUAL_INSTANTIATION
+#include "itkCuda2DSeparableConvolutionImageFilter.txx"
+#endif
+
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.txx CITK_CudaCanny//Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.txx
--- CITK_lastest_devel_build/Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.txx	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCuda2DSeparableConvolutionImageFilter.txx	2011-08-07 19:16:32.341186000 -0300
@@ -0,0 +1,137 @@
+/*=========================================================================
+
+Program:   Insight Segmentation & Registration Toolkit
+Module:    $RCSfile: itkCuda2DSeparableConvolutionImageFilter.txx,v $
+Language:  C++
+Date:      $Date: 2010-11-11 15:03:32 $
+Version:   $Revision: 0.1 $
+
+Copyright (c) Insight Software Consortium. All rights reser
+See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for detail.
+
+This software is distributed WITHOUT ANY WARRANTY; without even
+the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class Cuda2DSeparableConvolutionImageFilter
+ * \brief Class to convolve 2D images
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#ifndef __itkCuda2DSeparableConvolutionImageFilter_txx
+#define __itkCuda2DSeparableConvolutionImageFilter_txx
+
+#include "itkCuda2DSeparableConvolutionImageFilter.h"
+
+#include "itkImageBase.h"
+
+#define THREADS_PER_BLOCK 256 
+
+namespace itk {
+
+template<class TInputImage, class TOutputImage>
+Cuda2DSeparableConvolutionImageFilter<TInputImage, TOutputImage>
+::Cuda2DSeparableConvolutionImageFilter()
+{
+  m_CudaConf = CudaInterfaceType::New();
+}
+
+template<class TInputImage, class TOutputImage>
+Cuda2DSeparableConvolutionImageFilter<TInputImage, TOutputImage>
+::~Cuda2DSeparableConvolutionImageFilter()
+{
+}
+
+template<class TInputImage, class TOutputImage>
+void
+Cuda2DSeparableConvolutionImageFilter<TInputImage, TOutputImage>
+::GenerateData() 
+{
+
+  typename TInputImage::ConstPointer input = this->GetInput();
+  typename TOutputImage::Pointer output = this->GetOutput();
+  typename TOutputImage::PixelType * ptr;
+  typename TOutputImage::PixelType * d_Mask1;
+  typename TOutputImage::PixelType * d_Mask2;
+  
+  m_CudaConf->SetBlockDim(THREADS_PER_BLOCK,1,1);
+  m_CudaConf->SetGridDim((this->GetInput()->GetPixelContainer()->Size()+THREADS_PER_BLOCK-1)/THREADS_PER_BLOCK,1,1);
+
+  // Allocate output image object
+  output->SetBufferedRegion(output->GetRequestedRegion());
+
+  // Get image size
+  typename OutputImageType::SizeType size;
+  size = output->GetLargestPossibleRegion().GetSize();
+
+  // Copy the masks to the GPU memory. As they aren't image objects, this is not
+  // done by CITK
+  cudaMalloc((void**) &d_Mask1, (sizeof(typename TOutputImage::PixelType)*m_SizeMask1));
+  cudaMalloc((void**) &d_Mask2, (sizeof(typename TOutputImage::PixelType)*m_SizeMask2));
+  cudaMemcpy(d_Mask1, m_Mask1, (sizeof(typename TOutputImage::PixelType)*m_SizeMask1), cudaMemcpyHostToDevice);
+  cudaMemcpy(d_Mask2, m_Mask2, (sizeof(typename TOutputImage::PixelType)*m_SizeMask2), cudaMemcpyHostToDevice);
+
+  // Call cudaGaussian. Defined on CudaDiscreteGaussian.cu
+  ptr = cuda2DSeparableConvolution(m_CudaConf->GetGridDim(),m_CudaConf->GetBlockDim(),input->GetDevicePointer(), size[0], size[1], d_Mask1, m_SizeMask1, d_Mask2, m_SizeMask2);
+
+  // Set image pointer to the output image
+  output->GetPixelContainer()->SetDevicePointer(ptr, size[0]*size[1], true);
+
+  cudaFree(d_Mask1);
+  cudaFree(d_Mask2);
+  
+}
+
+
+/**
+ * Set Input Mask Horizontal
+ */
+template <typename TInputImage, typename TOutputImage>
+void
+Cuda2DSeparableConvolutionImageFilter<TInputImage,TOutputImage>
+::SetInputMaskHorizontal( OutputPixelType * input, unsigned int size )
+{
+
+  m_Mask1 = input;
+  m_SizeMask1 = size;
+}
+
+/**
+ * Get Input Mask Horizontal
+ */
+template <typename TInputImage, typename TOutputImage>
+const typename TOutputImage::PixelType *
+Cuda2DSeparableConvolutionImageFilter<TInputImage,TOutputImage>
+::GetInputMaskHorizontal( void )
+{
+  return dynamic_cast<const OutputPixelType *>( m_Mask1 );
+}
+
+
+/**
+ * Set Input Mask Vertical
+ */
+template <typename TInputImage, typename TOutputImage>
+void
+Cuda2DSeparableConvolutionImageFilter<TInputImage,TOutputImage>
+::SetInputMaskVertical( OutputPixelType * input, unsigned int size )
+{
+  m_Mask2 = input;
+  m_SizeMask2 = size;
+}
+
+/**
+ * Get Input Mask Vertical
+ */
+template <typename TInputImage, typename TOutputImage>
+const typename TOutputImage::PixelType *
+Cuda2DSeparableConvolutionImageFilter<TInputImage,TOutputImage>
+::GetInputMaskVertical( void )
+{
+  return dynamic_cast<const OutputPixelType *>( m_Mask2 );
+}
+
+}
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.h CITK_CudaCanny//Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.h
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.h	2011-08-07 19:16:51.989186001 -0300
@@ -0,0 +1,190 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCudaCannyEdgeDetectionImageFilter.h,v $
+  Language:  C++
+  Date:      $Date: 2010-09-15 12:27:15 $
+  Version:   $Revision: 3.0.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+#ifndef __itkCudaCannyEdgeDetectionImageFilter_h
+#define __itkCudaCannyEdgeDetectionImageFilter_h
+
+#include "itkImageToImageFilter.h"
+#include "itkImage.h"
+#include "itkCudaInterface.h"
+#include "itkCudaDiscreteGaussianImageFilter.h"
+#include "itkCudaZeroCrossingImageFilter.h"
+
+#include "cuda.h"
+#include "CudaCannyEdgeDetection.h"
+
+namespace itk
+{
+
+/** \class CannyEdgeDetectionImageFilter
+ *
+ * \brief This filter is an implementation of a Canny edge detector for scalar-valued
+ * images.  Based on John Canny's paper "A Computational Approach to Edge 
+ * Detection"(IEEE Transactions on Pattern Analysis and Machine Intelligence, 
+ * Vol. PAMI-8, No.6, November 1986),  there are four major steps used in the 
+ * edge-detection scheme:
+ * (1) Smooth the input image with Gaussian filter.
+ * (2) Calculate the second directional derivatives of the smoothed image. 
+ * (3) Non-Maximum Suppression: the zero-crossings of 2nd derivative are found,
+ *     and the sign of third derivative is used to find the correct extrema. 
+ * (4) The hysteresis thresholding is applied to the gradient magnitude
+ *      (multiplied with zero-crossings) of the smoothed image to find and 
+ *      link edges.
+ *
+ * \par Inputs and Outputs
+ * The input to this filter should be a scalar, real-valued Itk image of
+ * arbitrary dimension.  The output should also be a scalar, real-value Itk
+ * image of the same dimensionality.
+ *
+ * \par Parameters
+ * There are four parameters for this filter that control the sub-filters used
+ * by the algorithm.
+ *
+ * \par 
+ * Variance and Maximum error are used in the Gaussian smoothing of the input
+ * image.  See  itkDiscreteGaussianImageFilter for information on these
+ * parameters.
+ *
+ * \par
+ * Threshold is the lowest allowed value in the output image.  Its data type is 
+ * the same as the data type of the output image. Any values below the
+ * Threshold level will be replaced with the OutsideValue parameter value, whose
+ * default is zero.
+
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+
+template<class TInputImage, class TOutputImage>
+class ITK_EXPORT CudaCannyEdgeDetectionImageFilter
+  : public ImageToImageFilter<TInputImage, TOutputImage>
+{
+public:
+  /** Standard "Self" & Superclass typedef.  */
+  typedef CudaCannyEdgeDetectionImageFilter                 Self;
+  typedef ImageToImageFilter<TInputImage, TOutputImage> Superclass;
+   
+  /** Image typedef support   */
+  typedef TInputImage  InputImageType;
+  typedef TOutputImage OutputImageType;
+      
+  /** SmartPointer typedef support  */
+  typedef SmartPointer<Self>        Pointer;
+  typedef SmartPointer<const Self>  ConstPointer;
+
+  /** Define pixel types. */
+  typedef typename TInputImage::PixelType   InputImagePixelType;
+  typedef typename TOutputImage::PixelType  OutputImagePixelType;
+
+  /** Method for creation through the object factory. */
+  itkNewMacro(Self);  
+    
+  /** Typedef to describe the output image region type. */
+  typedef typename TOutputImage::RegionType OutputImageRegionType;
+  typedef typename TInputImage::RegionType  InputImageRegionType;
+    
+  /** Run-time type information (and related methods). */
+  itkTypeMacro(CudaCannyEdgeDetectionImageFilter, ImageToImageFilter);
+  
+  /** ImageDimension constant    */
+  itkStaticConstMacro(ImageDimension, unsigned int,
+                      TInputImage::ImageDimension);
+  itkStaticConstMacro(OutputImageDimension, unsigned int,
+                      TOutputImage::ImageDimension);
+ 
+   /** Typedef of double containers */
+  typedef FixedArray<double, itkGetStaticConstMacro(ImageDimension)> ArrayType;
+
+  /** Set the Variance parameter used by the Gaussian smoothing
+      filter in this algorithm */
+  virtual void SetVariance(const typename ArrayType::ValueType);
+
+ /** Set the MaximumError paramter used by the Gaussian smoothing filter
+      in this algorithm */
+  virtual void SetMaximumError(const typename ArrayType::ValueType);
+
+  ///* Set the Threshold value for detected edges. */
+  itkSetMacro(UpperThreshold, OutputImagePixelType );
+  itkGetConstMacro(UpperThreshold, OutputImagePixelType);
+  itkSetMacro(LowerThreshold, OutputImagePixelType );
+  itkGetConstMacro(LowerThreshold, OutputImagePixelType);
+
+#ifdef ITK_USE_CONCEPT_CHECKING
+  /** Begin concept checking */
+  itkConceptMacro(InputHasNumericTraitsCheck,
+    (Concept::HasNumericTraits<InputImagePixelType>));
+  itkConceptMacro(OutputHasNumericTraitsCheck,
+    (Concept::HasNumericTraits<OutputImagePixelType>));
+  itkConceptMacro(SameDimensionCheck,
+    (Concept::SameDimension<ImageDimension, OutputImageDimension>));
+  itkConceptMacro(InputIsFloatingPointCheck,
+    (Concept::IsFloatingPoint<InputImagePixelType>));
+  itkConceptMacro(OutputIsFloatingPointCheck,
+    (Concept::IsFloatingPoint<OutputImagePixelType>));
+  /** End concept checking */
+#endif
+
+protected:
+  CudaCannyEdgeDetectionImageFilter();
+  CudaCannyEdgeDetectionImageFilter(const Self&) {}
+  void PrintSelf(std::ostream& os, Indent indent) const;
+
+  void GenerateData();
+
+  typedef CudaDiscreteGaussianImageFilter<InputImageType, OutputImageType>
+                                                      CudaGaussianImageFilterType;
+
+  typedef CudaZeroCrossingImageFilter<OutputImageType, OutputImageType>
+                                                      CudaZeroCrossingFilterType;
+
+  typedef CudaInterface CudaInterfaceType;
+
+private:
+  virtual ~CudaCannyEdgeDetectionImageFilter(){};
+
+  /** Implementation of 2md Derivative and Magnitude calculation on Cuda*/
+  void Cuda2ndDerivative();
+
+  /** Implementation of Hysteresis Thresholding on Cuda */
+  void CudaHysteresisThresholding();
+
+  /** Upper threshold value for identifying edges. */
+  OutputImagePixelType m_UpperThreshold;
+  
+  /** Lower threshold value for identifying edges. */
+  OutputImagePixelType m_LowerThreshold;
+
+  /** Update buffers used during calculationsof multiple steps */
+  typename OutputImageType::Pointer m_UpdateBuffer1;
+
+  /** CudaGaussian filter to smooth the input image  */
+  typename CudaGaussianImageFilterType::Pointer m_CudaGaussianFilter;
+ 
+  /** CudaZeroCrossing filter to detect zero crossings on the 2nd derivative
+   * image */
+  typename CudaZeroCrossingFilterType::Pointer m_CudaZeroCrossingFilter;
+
+  typename CudaInterfaceType::Pointer m_CudaConf;
+};
+
+} //end of namespace itk
+
+#ifndef ITK_MANUAL_INSTANTIATION
+#include "itkCudaCannyEdgeDetectionImageFilter.txx"
+#endif
+  
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.txx CITK_CudaCanny//Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.txx
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.txx	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaCannyEdgeDetectionImageFilter.txx	2011-08-07 19:16:52.065185999 -0300
@@ -0,0 +1,208 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCudaCannyEdgeDetectionImageFilter.txx,v $
+  Language:  C++
+  Date:      $Date: 2010-09-15 12:01:33 $
+  Version:   $Revision: 3.0.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class CannyEdgeDetectionImageFilter
+ * \brief CUDA implementation of Canny's algorithm
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+#ifndef __itkCudaCannyEdgeDetectionImageFilter_txx
+#define __itkCudaCannyEdgeDetectionImageFilter_txx
+#include "itkCudaCannyEdgeDetectionImageFilter.h"
+
+#define THREADS_PER_BLOCK 256 
+
+namespace itk
+{
+  
+template <class TInputImage, class TOutputImage>
+CudaCannyEdgeDetectionImageFilter<TInputImage, TOutputImage>::
+CudaCannyEdgeDetectionImageFilter()
+{
+
+  m_UpperThreshold = NumericTraits<OutputImagePixelType>::Zero;
+  m_LowerThreshold = NumericTraits<OutputImagePixelType>::Zero;
+
+  m_CudaGaussianFilter = CudaGaussianImageFilterType::New();
+  m_CudaZeroCrossingFilter = CudaZeroCrossingFilterType::New();
+  m_CudaConf = CudaInterfaceType::New();
+
+  m_UpdateBuffer1  = OutputImageType::New();
+
+}
+
+template <class TInputImage, class TOutputImage>
+void 
+CudaCannyEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::SetVariance(const typename ArrayType::ValueType v)
+{
+  // call the class' implementation of this method
+  m_CudaGaussianFilter->SetVariance(v);
+  return;
+}
+
+template <class TInputImage, class TOutputImage>
+void 
+CudaCannyEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::SetMaximumError(const typename ArrayType::ValueType v)
+{
+  // call the class' implementation of this method
+  m_CudaGaussianFilter->SetMaximumError(v);
+  return;
+}
+
+template <class TInputImage, class TOutputImage>
+void 
+CudaCannyEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::Cuda2ndDerivative()
+{
+
+  // Set input, output and temporary pointers
+  typename OutputImageType::Pointer input  = m_CudaGaussianFilter->GetOutput();
+  typename OutputImageType::Pointer output = this->GetOutput();
+  typename OutputImageType::Pointer output1 = m_UpdateBuffer1;
+  typename TOutputImage::PixelType * deriv;
+  typename TOutputImage::PixelType * mag;
+
+  // Get image size
+  typename OutputImageType::SizeType size = output->GetLargestPossibleRegion().GetSize();
+
+  // Call CudaNMS. Defined on CudaCannyEdgeDetection.cu
+  deriv = cuda2ndDerivative(m_CudaConf->GetGridDim(),m_CudaConf->GetBlockDim(),input->GetDevicePointer(),size[0],size[1]);
+
+  mag = cuda2ndDerivativePos(m_CudaConf->GetGridDim(),m_CudaConf->GetBlockDim(),input->GetDevicePointer(),deriv,size[0],size[1]);
+
+  // Set NMS pointer on the output image
+  output->GetPixelContainer()->SetDevicePointer(deriv, size[0]*size[1], true);
+  output1->GetPixelContainer()->SetDevicePointer(mag, size[0]*size[1], true);
+
+}
+
+template <class TInputImage, class TOutputImage>
+void 
+CudaCannyEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::CudaHysteresisThresholding()
+{
+
+  // Get the image ponter
+  typename OutputImageType::Pointer mag    = m_UpdateBuffer1;
+  typename OutputImageType::Pointer deriv  = m_CudaZeroCrossingFilter->GetOutput();
+  typename OutputImageType::Pointer output = this->GetOutput();
+  typename TOutputImage::PixelType * edges;
+
+  // Get image size
+  typename OutputImageType::SizeType size = output->GetLargestPossibleRegion().GetSize();
+
+  // Call CudaHysteresis. Defined on CudaCannyEdgeDetection.cu
+  // This Multiplyes the Zero crossings of the Second derivative with the
+  // magnitude gradients of the image.
+  edges = cudaHysteresis(m_CudaConf->GetGridDim(),m_CudaConf->GetBlockDim(),deriv->GetDevicePointer(), mag->GetDevicePointer(), size[0], size[1], this->GetLowerThreshold(), this->GetUpperThreshold());
+
+  // Set Canny output
+  output->GetPixelContainer()->SetDevicePointer(edges, size[0]*size[1], true);
+}
+
+template< class TInputImage, class TOutputImage >
+void
+CudaCannyEdgeDetectionImageFilter< TInputImage, TOutputImage >
+::GenerateData()
+{
+
+  // Get input and output image pointers
+  typename InputImageType::ConstPointer input = this->GetInput();
+  typename OutputImageType::Pointer output = this->GetOutput();
+ 
+  // Allocate output image object
+  m_UpdateBuffer1->CopyInformation( input );
+  m_UpdateBuffer1->SetRequestedRegion(input->GetRequestedRegion());
+  m_UpdateBuffer1->SetBufferedRegion(input->GetBufferedRegion());
+  m_UpdateBuffer1->Allocate();  
+  output->SetBufferedRegion(output->GetRequestedRegion());
+  
+  m_CudaConf->SetBlockDim(THREADS_PER_BLOCK,1,1);
+  m_CudaConf->SetGridDim((this->GetInput()->GetPixelContainer()->Size()+THREADS_PER_BLOCK-1)/THREADS_PER_BLOCK,1,1);
+
+  // 1.Apply the Gaussian Filter to the input image.-------
+  m_CudaGaussianFilter->SetInput(input);
+  
+  unsigned int timer = 0;
+  cutCreateTimer( &timer );
+  cutStartTimer( timer );  /// Start timer
+
+  m_CudaGaussianFilter->Update();
+
+  cutStopTimer( timer );  /// Stop timer
+  printf("Gaussian time = %f ms\n",cutGetTimerValue( timer ));
+
+  // 2. Calculate 2nd order directional derivative.-------
+  // Calculate the 2nd order directional derivative of the smoothed image.
+  // The output of this filter will be used to store the directional
+  // derivative.
+  timer = 0;
+  cutCreateTimer( &timer );
+  cutStartTimer( timer );  ///< Start timer
+
+  this->Cuda2ndDerivative();
+
+  cutStopTimer( timer );  ///< Stop timer
+  printf("2nd Derivative time = %f ms\n",cutGetTimerValue( timer ));
+
+  // 3. NonMaximumSupression.-------
+  // Calculate the zero crossings of the 2nd directional derivative and write 
+  // the result to output buffer. 
+  m_CudaZeroCrossingFilter->SetInput(this->GetOutput());
+
+  timer = 0;
+  cutCreateTimer( &timer );
+  cutStartTimer( timer );  ///< Start timer
+
+  m_CudaZeroCrossingFilter->Update();
+
+  cutStopTimer( timer );  ///< Stop timer
+  printf("Maximum Supression time = %f ms\n",cutGetTimerValue( timer ));
+
+  // 4. Hysteresis Thresholding.-------
+  // Get all the edges corresponding to zerocrossings by multiplying the 2nd
+  // directional derivative and the magnitude gradient on preparation.
+  // Then do the double thresholding upon the edge responses.
+  timer = 0;
+  cutCreateTimer( &timer );
+  cutStartTimer( timer );  ///< Start timer
+
+  this->CudaHysteresisThresholding();
+
+  cutStopTimer( timer );  ///< Stop timer
+  printf("Hysteresis time = %f ms\n",cutGetTimerValue( timer ));
+
+}
+
+template <class TInputImage, class TOutputImage>
+void 
+CudaCannyEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::PrintSelf(std::ostream& os, Indent indent) const
+{
+  Superclass::PrintSelf(os,indent);
+
+  os << indent << "UpperThreshold: "
+     << static_cast<typename NumericTraits<OutputImagePixelType>::PrintType>(m_UpperThreshold)
+     << std::endl;
+  os << indent << "LowerThreshold: "
+     << static_cast<typename NumericTraits<OutputImagePixelType>::PrintType>(m_LowerThreshold)
+     << std::endl;
+}
+
+}//end of itk namespace
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.h CITK_CudaCanny//Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.h
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.h	2011-08-07 19:17:13.269186000 -0300
@@ -0,0 +1,247 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCudaDiscreteGaussianImageFilter.h,v $
+  Language:  C++
+  Date:      $Date: 2010-11-10 12:27:21 $
+  Version:   $Revision: 1.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+#ifndef __itkCudaDiscreteGaussianImageFilter_h
+#define __itkCudaDiscreteGaussianImageFilter_h
+
+#include "itkImageToImageFilter.h"
+#include "itkCudaInterface.h"
+#include "itkCuda2DSeparableConvolutionImageFilter.h"
+#include "itkGaussianOperator.h"
+#include "itkImage.h"
+
+#include "cuda.h"
+#include "CudaDiscreteGaussian.h"
+
+namespace itk
+{
+/**
+ * \class CudaDiscreteGaussianImageFilter
+ * \brief Blurs an image by separable convolution with discrete gaussian kernels.
+ * This filter performs Gaussian blurring by separable convolution of an image
+ * and a discrete Gaussian operator (kernel).
+ *
+ * \sa cuda1DGaussianOperator
+ * \sa Image
+ * 
+ * \ingroup ImageEnhancement 
+ * \ingroup ImageFeatureExtraction 
+ *
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+ */
+
+template <class TInputImage, class TOutputImage >
+class ITK_EXPORT CudaDiscreteGaussianImageFilter :
+    public ImageToImageFilter< TInputImage, TOutputImage > 
+{
+public:
+  /** Standard class typedefs. */
+  typedef CudaDiscreteGaussianImageFilter                 Self;
+  typedef ImageToImageFilter< TInputImage, TOutputImage > Superclass;
+  typedef SmartPointer<Self>                              Pointer;
+  typedef SmartPointer<const Self>                        ConstPointer;
+
+  /** Method for creation through the object factory. */
+  itkNewMacro(Self);
+
+  /** Run-time type information (and related methods). */
+  itkTypeMacro(CudaDiscreteGaussianImageFilter, ImageToImageFilter);
+  
+  /** Image type information. */
+  typedef TInputImage  InputImageType;
+  typedef TOutputImage OutputImageType;
+
+  /** Extract some information from the image types.  Dimensionality
+   * of the two images is assumed to be the same. */
+  typedef typename TOutputImage::PixelType         OutputPixelType;
+  typedef typename TOutputImage::InternalPixelType OutputInternalPixelType;
+  typedef typename TInputImage::PixelType          InputPixelType;
+  typedef typename TInputImage::InternalPixelType  InputInternalPixelType;
+  typedef Image<OutputPixelType,1> MaskImageType;
+
+  TOutputImage * GetOutput();
+
+  /** Extract some information from the image types.  Dimensionality
+   * of the two images is assumed to be the same. */
+  itkStaticConstMacro(ImageDimension, unsigned int,
+                      TOutputImage::ImageDimension);
+  
+  /** Typedef of double containers */
+  typedef FixedArray<double, itkGetStaticConstMacro(ImageDimension)> ArrayType;
+
+  /** The variance for the discrete Gaussian kernel.  Sets the variance
+   * independently for each dimension, but 
+   * see also SetVariance(const double v). The default is 0.0 in each
+   * dimension. If UseImageSpacing is true, the units are the physical units
+   * of your image.  If UseImageSpacing is false then the units are
+   * pixels. */
+  itkSetMacro(Variance, ArrayType);
+  itkGetConstMacro(Variance, const ArrayType);
+
+  /** The algorithm will size the discrete kernel so that the error
+   * resulting from truncation of the kernel is no greater than
+   * MaximumError. The default is 0.01 in each dimension. */
+  itkSetMacro(MaximumError, ArrayType);
+  itkGetConstMacro(MaximumError, const ArrayType);
+
+
+  /** Set the kernel to be no wider than MaximumKernelWidth pixels,
+   *  even if MaximumError demands it. The default is 3 pixels. */
+  itkGetConstMacro(MaximumKernelWidth, unsigned int);
+  itkSetMacro(MaximumKernelWidth, unsigned int);
+
+  /** Set the number of dimensions to smooth. Defaults to the image
+   * dimension. Can be set to less than ImageDimension, smoothing all
+   * the dimensions less than FilterDimensionality.  For instance, to
+   * smooth the slices of a volume without smoothing in Z, set the
+   * FilterDimensionality to 2. */
+  itkGetConstMacro(FilterDimensionality, unsigned int);
+  itkSetMacro(FilterDimensionality, unsigned int);
+
+  /** Convenience Set methods for setting all dimensional parameters
+   *  to the same values. */
+  void SetVariance (const typename ArrayType::ValueType v)
+    {
+    m_Variance.Fill(v);
+    }
+
+  void SetMaximumError (const typename ArrayType::ValueType v)
+    {
+    m_MaximumError.Fill(v);
+    }
+
+  void SetVariance (const double *v)
+    {
+    ArrayType dv;
+    for (unsigned int i = 0; i < ImageDimension; i++)
+      {
+      dv[i] = v[i];
+      }
+    this->SetVariance(dv);
+    }
+
+  void SetVariance (const float *v)
+    {
+    ArrayType dv;
+    for (unsigned int i = 0; i < ImageDimension; i++)
+      {
+      dv[i] = v[i];
+      }
+    this->SetVariance(dv);
+    }
+
+  void SetMaximumError (const double *v)
+    {
+    ArrayType dv;
+    for (unsigned int i = 0; i < ImageDimension; i++)
+      {
+      dv[i] = v[i];
+      }
+    this->SetMaximumError(dv);
+    }
+
+  void SetMaximumError (const float *v)
+    {
+    ArrayType dv;
+    for (unsigned int i = 0; i < ImageDimension; i++)
+      {
+      dv[i] = v[i];
+      }
+    this->SetMaximumError(dv);
+    }
+
+  /** Use the image spacing information in calculations. Use this option if you
+   *  want to specify Gaussian variance in real world units.  Default is
+   *   ImageSpacingOn. */
+  void SetUseImageSpacingOn()
+    { this->SetUseImageSpacing(true); }
+  
+  /** Ignore the image spacing. Use this option if you want to specify Gaussian
+      variance in pixels.  Default is ImageSpacingOn. */
+  void SetUseImageSpacingOff()
+    { this->SetUseImageSpacing(false); }
+  
+  /** Set/Get whether or not the filter will use the spacing of the input
+      image in its calculations */
+  itkSetMacro(UseImageSpacing, bool);
+  itkGetConstMacro(UseImageSpacing, bool);
+
+#ifdef ITK_USE_CONCEPT_CHECKING
+  /** Begin concept checking */
+  itkConceptMacro(OutputHasNumericTraitsCheck,
+    (Concept::HasNumericTraits<OutputPixelType>));
+  /** End concept checking */
+#endif
+
+protected:
+  CudaDiscreteGaussianImageFilter()
+    {
+    m_Variance.Fill(0.0);
+    m_MaximumError.Fill(0.01);
+    m_MaximumKernelWidth = 32;
+    m_UseImageSpacing = false;
+    m_FilterDimensionality = ImageDimension;
+    m_CudaConvolutionFilter = Cuda2DSeparableConvolutionImageFilterType::New();
+    m_CudaConf = CudaInterfaceType::New();
+    }
+  virtual ~CudaDiscreteGaussianImageFilter() {}
+  void PrintSelf(std::ostream& os, Indent indent) const;
+
+  /** Standard pipeline method. */
+  void GenerateData();
+
+  typedef Cuda2DSeparableConvolutionImageFilter<InputImageType, OutputImageType>
+                                                      Cuda2DSeparableConvolutionImageFilterType;
+  typedef CudaInterface CudaInterfaceType;
+  
+private:
+  CudaDiscreteGaussianImageFilter(const Self&); //purposely not implemented
+  void operator=(const Self&); //purposely not implemented
+
+  /** The variance of the gaussian blurring kernel in each dimensional direction. */
+  ArrayType m_Variance;
+
+  /** The maximum error of the gaussian blurring kernel in each dimensional
+   * direction. For definition of maximum error, see GaussianOperator.
+   * \sa GaussianOperator */
+  ArrayType m_MaximumError;
+
+  /** Maximum allowed kernel width for any dimension of the discrete Gaussian
+      approximation */
+  unsigned int m_MaximumKernelWidth;
+
+  /** Number of dimensions to process. Default is all dimensions */
+  unsigned int m_FilterDimensionality;
+
+  /** Flag to indicate whether to use image spacing */
+  bool m_UseImageSpacing;
+
+  /** Cuda2DSeparableConvolution to convolve the image  */
+  typename Cuda2DSeparableConvolutionImageFilterType::Pointer m_CudaConvolutionFilter;
+
+  typename CudaInterfaceType::Pointer m_CudaConf;
+  
+  GaussianOperator<OutputPixelType, ImageDimension> oper;
+};
+  
+} // end namespace itk
+
+#ifndef ITK_MANUAL_INSTANTIATION
+#include "itkCudaDiscreteGaussianImageFilter.txx"
+#endif
+
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.txx CITK_CudaCanny//Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.txx
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.txx	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaDiscreteGaussianImageFilter.txx	2011-08-07 19:17:13.341186000 -0300
@@ -0,0 +1,143 @@
+/*=========================================================================
+
+Program:   Insight Segmentation & Registration Toolkit
+Module:    $RCSfile: itkCudaDiscreteGaussianImageFilter.txx,v $
+Language:  C++
+Date:      $Date: 2010-11-10 12:44:26 $
+Version:   $Revision: 1.0 $
+
+Copyright (c) Insight Software Consortium. All rights reserved.
+See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+This software is distributed WITHOUT ANY WARRANTY; without even 
+the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/**
+ * \class CudaDiscreteGaussianImageFilter
+ * \brief Blurs an image by separable convolution with discrete gaussian kernels.
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#ifndef __itkCudaDiscreteGaussianImageFilter_txx
+#define __itkCudaDiscreteGaussianImageFilter_txx
+
+#include "itkCudaDiscreteGaussianImageFilter.h"
+#include "itkGaussianOperator.h"
+
+#define THREADS_PER_BLOCK 256 
+
+namespace itk
+{
+
+template< class TInputImage, class TOutputImage >
+void
+CudaDiscreteGaussianImageFilter<TInputImage, TOutputImage>
+::GenerateData()
+{
+  typename TInputImage::ConstPointer input = this->GetInput();
+  typename TOutputImage::Pointer output = this->GetOutput();
+//  typename TOutputImage::PixelType * mask;
+
+  // Determine the dimensionality to filter
+  unsigned int filterDimensionality = m_FilterDimensionality;
+  if (filterDimensionality > ImageDimension)
+    {
+    filterDimensionality = ImageDimension;
+    }
+  if (filterDimensionality == 0)
+    {
+    // no smoothing, copy input to output
+    ImageRegionConstIterator<InputImageType> inIt(
+      input,
+      this->GetOutput()->GetRequestedRegion() );
+    ImageRegionIterator<OutputImageType> outIt(
+      output,
+      this->GetOutput()->GetRequestedRegion() );
+
+    while (!inIt.IsAtEnd())
+      {
+      outIt.Set( static_cast<OutputPixelType>(inIt.Get()) );
+      ++inIt;
+      ++outIt;
+      }
+    return;
+    }
+
+   // Create a series of operators
+  typedef GaussianOperator<OutputPixelType, ImageDimension> OperatorType;
+  std::vector<OperatorType> oper;
+  oper.resize(filterDimensionality);
+
+  // Set up the operators
+  unsigned int i;
+  for (i = 0; i < filterDimensionality; ++i){
+    // we reverse the direction to minimize computation while, because
+    // the largest dimension will be split slice wise for streaming 
+    unsigned int reverse_i = filterDimensionality - i - 1;
+
+    // Set up the operator for this dimension
+    oper[reverse_i].SetDirection(i);
+    if (m_UseImageSpacing == true)
+      {
+      if (input->GetSpacing()[i] == 0.0)
+        {
+        itkExceptionMacro(<< "Pixel spacing cannot be zero");
+        }
+      else
+        {
+        // convert the variance from physical units to pixels
+        double s = input->GetSpacing()[i];
+        s = s*s;
+        oper[reverse_i].SetVariance(m_Variance[i] / s);
+        }
+      }
+    else
+      {
+      oper[reverse_i].SetVariance(m_Variance[i]);
+      }
+
+    oper[reverse_i].SetMaximumKernelWidth(m_MaximumKernelWidth);
+    oper[reverse_i].SetMaximumError(m_MaximumError[i]);
+    oper[reverse_i].CreateDirectional();
+  }
+
+  m_CudaConf->SetBlockDim(oper[0].Size(),1,1);
+  m_CudaConf->SetGridDim(1,1,1);
+
+//  mask = cuda1DGaussianOperator(m_CudaConf->GetGridDim(),m_CudaConf->GetBlockDim(),this->GetMaximumKernelWidth(),(float) this->GetVariance());
+
+  m_CudaConvolutionFilter->SetInput(input);
+  m_CudaConvolutionFilter->SetInputMaskHorizontal(oper[0].GetBufferReference().begin(), oper[0].Size());
+  m_CudaConvolutionFilter->SetInputMaskVertical(oper[1].GetBufferReference().begin(), oper[1].Size());
+  m_CudaConvolutionFilter->Update();
+
+}
+
+template< class TInputImage, class TOutputImage >
+TOutputImage *
+CudaDiscreteGaussianImageFilter<TInputImage, TOutputImage>
+::GetOutput()
+{
+
+  return(m_CudaConvolutionFilter->GetOutput());
+
+}
+
+template< class TInputImage, class TOutputImage >
+void
+CudaDiscreteGaussianImageFilter<TInputImage, TOutputImage>
+::PrintSelf(std::ostream& os, Indent indent) const
+{
+  Superclass::PrintSelf(os,indent);
+
+  os << indent << "Variance: " << m_Variance << std::endl;
+  os << indent << "MaximumKernelWidth: " << m_MaximumKernelWidth << std::endl;
+  os << indent << "FilterDimensionality: " << m_FilterDimensionality << std::endl;
+}
+
+} // end namespace itk
+
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.h CITK_CudaCanny//Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.h
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.h	2011-08-07 19:17:31.925186001 -0300
@@ -0,0 +1,156 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCudaSobelEdgeDetectionImageFilter.h,v $
+  Language:  C++
+  Date:      $Date: 2010-09-15 20:49:56 $
+  Version:   $Revision: 3.0.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+#ifndef __itkCudaSobelEdgeDetectionImageFilter_h
+#define __itkCudaSobelEdgeDetectionImageFilter_h
+
+#include "itkImageToImageFilter.h"
+#include "itkImage.h"
+
+#include "cuda.h"
+#include "CudaSobelEdgeDetection.h"
+
+namespace itk
+{
+
+/**
+ * \class SobelEdgeDetectionImageFilter
+ * \brief A 2D edge detection using the Sobel operator implemented in CUDA.
+ *
+ * This filter uses the Sobel operator to calculate the image gradient and then
+ * finds the magnitude and direction of this gradient vector.  The Sobel gradient 
+ * magnitude (square-root sum of squares) is an indication of edge strength.
+ * 
+ * \sa ImageToImageFilter
+ * 
+ * \ingroup ImageFeatureExtraction 
+ *
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+ */
+
+template <class TInputImage, class TOutputImage>
+class ITK_EXPORT CudaSobelEdgeDetectionImageFilter : 
+    public ImageToImageFilter< TInputImage, TOutputImage > 
+{
+public:
+  /**
+   * Standard "Self" & Superclass typedef.
+   */
+  typedef CudaSobelEdgeDetectionImageFilter                   Self;
+  typedef ImageToImageFilter< TInputImage, TOutputImage > Superclass;
+
+  /**
+   * Extract some information from the image types.  Dimensionality
+   * of the two images is assumed to be the same.
+   */
+  typedef typename TOutputImage::PixelType         OutputPixelType;
+  typedef typename TOutputImage::InternalPixelType OutputInternalPixelType;
+  typedef typename TInputImage::PixelType          InputPixelType;
+  typedef typename TInputImage::InternalPixelType  InputInternalPixelType;
+  typedef float                                    DirectionPixelType;
+  typedef itk::Image<DirectionPixelType,2>         DirectionType;
+  itkStaticConstMacro(ImageDimension, unsigned int,
+                      TOutputImage::ImageDimension );
+  itkStaticConstMacro(InputImageDimension, unsigned int,
+                      TInputImage::ImageDimension );
+  
+  /**
+   * Image typedef support
+   */
+  typedef TInputImage                      InputImageType;
+  typedef TOutputImage                     OutputImageType;
+  typedef typename InputImageType::Pointer InputImagePointer;
+
+  /** 
+   * Smart pointer typedef support 
+   */
+  typedef SmartPointer<Self>        Pointer;
+  typedef SmartPointer<const Self>  ConstPointer;
+  
+  /**
+   * Run-time type information (and related methods)
+   */
+  itkTypeMacro(CudaSobelEdgeDetectionImageFilter, ImageToImageFilter);
+  
+  /**
+   * Method for creation through the object factory.
+   */
+  itkNewMacro(Self);
+
+  /** Method for getting the gradient direction of each pixel */
+  DirectionType* GetGradientDirection();
+
+  /** Method for getting the gradient strenght/magnitude of each pixel */
+  OutputImageType* GetGradientMagnitude();
+  
+  /**
+   * SobelEdgeDetectionImageFilter needs a larger input requested region than
+   * the output requested region (larger in the direction of the
+   * derivative).  As such, SobelEdgeDetectionImageFilter needs to provide an
+   * implementation for GenerateInputRequestedRegion() in order to
+   * inform the pipeline execution model.
+   *
+   * \sa ImageToImageFilter::GenerateInputRequestedRegion()
+   */
+  virtual void GenerateInputRequestedRegion() throw(InvalidRequestedRegionError);
+
+#ifdef ITK_USE_CONCEPT_CHECKING
+  /** Begin concept checking */
+  itkConceptMacro(SameDimensionCheck,
+    (Concept::SameDimension<InputImageDimension, ImageDimension>));
+  itkConceptMacro(OutputHasNumericTraitsCheck,
+    (Concept::HasNumericTraits<OutputPixelType>));
+
+#ifdef ITK_USE_STRICT_CONCEPT_CHECKING
+  itkConceptMacro(OutputPixelIsFloatingPointCheck,
+    (Concept::IsFloatingPoint<OutputPixelType>));
+#endif
+
+  /** End concept checking */
+#endif
+
+protected:
+  CudaSobelEdgeDetectionImageFilter();
+  virtual ~CudaSobelEdgeDetectionImageFilter() {}
+  CudaSobelEdgeDetectionImageFilter(const Self&) {}
+  //void operator=(const Self&) {}
+
+  /**
+   * Standard pipeline method. This class does not implement a
+   * ThreadedGenerateData(), its GenerateData() call an CUDA function to
+   * paralelize the operation.
+   */
+  void GenerateData();
+  void PrintSelf(std::ostream& os, Indent indent) const
+  { Superclass::PrintSelf(os,indent); }
+
+  
+private:
+
+  /** GPU structure to the Gradient of the Image. 
+   * Tgrad is defined on CudaSobelEdgeDetection.cu */
+  Tgrad m_gradient;
+
+};
+  
+} // end namespace itk
+
+#ifndef ITK_MANUAL_INSTANTIATION
+#include "itkCudaSobelEdgeDetectionImageFilter.txx"
+#endif
+
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.txx CITK_CudaCanny//Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.txx
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.txx	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaSobelEdgeDetectionImageFilter.txx	2011-08-07 19:17:31.977186001 -0300
@@ -0,0 +1,103 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCudaSobelEdgeDetectionImageFilter.txx,v $
+  Language:  C++
+  Date:      $Date: 2010-09-15 20:50:03 $
+  Version:   $Revision: 3.0.0 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/**
+ * \class SobelEdgeDetectionImageFilter
+ * \brief A 2D edge detection using the Sobel operator implemented in CUDA.
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#ifndef __itkCudaSobelEdgeDetectionImageFilter_txx
+#define __itkCudaSobelEdgeDetectionImageFilter_txx
+#include "itkCudaSobelEdgeDetectionImageFilter.h"
+
+
+namespace itk
+{
+
+template <class TInputImage, class TOutputImage> 
+CudaSobelEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::CudaSobelEdgeDetectionImageFilter()
+{
+ 
+  typename TOutputImage::Pointer magnitude   = OutputImageType::New();
+  typename DirectionType::Pointer direction = DirectionType::New();
+
+  this->SetNumberOfRequiredOutputs( 2 );
+  this->SetNthOutput( 0, magnitude.GetPointer() );
+  this->SetNthOutput( 1, direction.GetPointer() );
+
+}
+
+template <class TInputImage, class TOutputImage>
+void 
+CudaSobelEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::GenerateInputRequestedRegion() throw (InvalidRequestedRegionError)
+{
+  // call the superclass' implementation of this method. this should
+  // copy the output requested region to the input requested region
+  Superclass::GenerateInputRequestedRegion();
+  
+}
+
+template <class TInputImage, class TOutputImage>
+typename CudaSobelEdgeDetectionImageFilter<TInputImage,TOutputImage>::DirectionType * 
+CudaSobelEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::GetGradientDirection()
+{
+  return( dynamic_cast<DirectionType *>(this->ProcessObject::GetOutput( 1 )) );
+}
+
+template <class TInputImage, class TOutputImage>
+TOutputImage *
+CudaSobelEdgeDetectionImageFilter<TInputImage,TOutputImage>
+::GetGradientMagnitude()
+{
+  return( dynamic_cast<TOutputImage *>(this->ProcessObject::GetOutput( 0 )) );
+}
+
+template< class TInputImage, class TOutputImage >
+void
+CudaSobelEdgeDetectionImageFilter< TInputImage, TOutputImage >
+::GenerateData()
+{
+
+  // Set input, output and temporary pointers
+  typename TInputImage::ConstPointer input  = this->GetInput();
+  typename TOutputImage::Pointer magnitude  = this->GetGradientMagnitude();
+  typename DirectionType::Pointer direction = this->GetGradientDirection();
+  Tgrad * tmpGrad;
+  
+  // Allocate output image object
+  magnitude->SetBufferedRegion(magnitude->GetRequestedRegion());
+  direction->SetBufferedRegion(direction->GetRequestedRegion());
+
+  // Get image size
+  typename OutputImageType::SizeType size = magnitude->GetLargestPossibleRegion().GetSize();
+
+  // Call cudaSobel. Defined on CudaSobelEdgeDetection.cu
+  tmpGrad = cudaSobel(&m_gradient, input->GetDevicePointer(), size[0], size[1]);
+
+  // Set image pointer to the output image
+  magnitude->GetPixelContainer()->SetDevicePointer(tmpGrad->Magnitude, size[0]*size[1], true);
+  direction->GetPixelContainer()->SetDevicePointer(tmpGrad->Direction, size[0]*size[1], true);
+
+}
+
+} // end namespace itk
+
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaZeroCrossingImageFilter.h CITK_CudaCanny//Code/BasicFilters/itkCudaZeroCrossingImageFilter.h
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaZeroCrossingImageFilter.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaZeroCrossingImageFilter.h	2011-08-07 19:17:50.293186001 -0300
@@ -0,0 +1,148 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCudaZeroCrossingImageFilter.h,v $
+  Language:  C++
+  Date:      $Date: 2010-11-04 12:28:13 $
+  Version:   $Revision: 0.1 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+#ifndef __itkCudaZeroCrossingImageFilter_h
+#define __itkCudaZeroCrossingImageFilter_h
+
+#include "itkImageToImageFilter.h"
+#include "itkImage.h"
+#include "itkCudaInterface.h"
+
+#include "cuda.h"
+#include "CudaZeroCrossing.h"
+
+namespace itk
+{  
+/** \class CudaZeroCrossingImageFilter
+ *
+ *  This filter finds the closest pixel to the zero-crossings (sign changes) in
+ *  a signed itk::Image.  Pixels closest to zero-crossings are labeled with
+ *  a foreground value.  All other pixels are marked with a background value.
+ *  The algorithm works by detecting differences in sign among neighbors using
+ *  city-block style connectivity (4-neighbors in 2d, 6-neighbors in 3d, etc.).
+ *  
+ *  \par Inputs and Outputs
+ *  The input to this filter is an itk::Image of arbitrary dimension.  The
+ *  algorithm assumes a signed data type (zero-crossings are not defined for
+ *  unsigned data types), and requires that operator>, operator<, operator==,
+ *  and operator!= are defined.  
+ *
+ *  \par
+ *  The output of the filter is a binary, labeled image of user-specified type.
+ *  By default, zero-crossing pixels are labeled with a default ``foreground''
+ *  value of itk::NumericTraits<OutputDataType>::One, where OutputDataType is
+ *  the data type of the output image.  All other pixels are labeled with a
+ *  default ``background'' value of itk::NumericTraits<OutputDataType>::Zero.
+ *
+ *  \par Parameters
+ *  There are two parameters for this filter.  ForegroundValue is the value
+ *  that marks zero-crossing pixels.  The BackgroundValue is the value given to 
+ *  all other pixels.
+ *
+ *  \sa Image
+ *  \ingroup ImageFeatureExtraction 
+ *
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+template<class TInputImage, class TOutputImage>
+class ITK_EXPORT CudaZeroCrossingImageFilter
+  : public ImageToImageFilter<TInputImage, TOutputImage>
+{
+public:
+  /** Standard "Self" & Superclass typedef. */
+  typedef CudaZeroCrossingImageFilter                   Self;
+  typedef ImageToImageFilter<TInputImage, TOutputImage> Superclass;
+  
+  /** Image typedef support   */
+  typedef TInputImage  InputImageType;
+  typedef TOutputImage OutputImageType;
+  
+  /** SmartPointer typedef support  */ 
+  typedef SmartPointer<Self>        Pointer;
+  typedef SmartPointer<const Self>  ConstPointer;
+  
+  /** Define pixel types  */
+  typedef typename TInputImage::PixelType   InputImagePixelType;
+  typedef typename TOutputImage::PixelType  OutputImagePixelType;
+  
+  /** Method for creation through the object factory.  */
+  itkNewMacro(Self);  
+  
+  /** Typedef to describe the output image region type. */
+  typedef typename TOutputImage::RegionType OutputImageRegionType;
+  
+  /** Run-time type information (and related methods). */
+  itkTypeMacro(CudaZeroCrossingImageFilter, ImageToImageFilter);
+  
+  /** ImageDimension enumeration   */
+  itkStaticConstMacro(ImageDimension, unsigned int,
+                      TInputImage::ImageDimension );
+  itkStaticConstMacro(OutputImageDimension, unsigned int,
+                      TOutputImage::ImageDimension );
+  
+  /** Set/Get the label value for zero-crossing pixels. */
+  itkSetMacro(ForegroundValue, OutputImagePixelType);
+  itkGetConstMacro(ForegroundValue, OutputImagePixelType);
+  
+  /** Set/Get the label value for non-zero-crossing pixels. */
+  itkSetMacro(BackgroundValue, OutputImagePixelType);
+  itkGetConstMacro(BackgroundValue, OutputImagePixelType);
+
+#ifdef ITK_USE_CONCEPT_CHECKING
+  /** Begin concept checking */
+  itkConceptMacro(OutputEqualityComparableCheck,
+    (Concept::EqualityComparable<OutputImagePixelType>));
+  itkConceptMacro(SameDimensionCheck,
+    (Concept::SameDimension<ImageDimension, OutputImageDimension>));
+  itkConceptMacro(InputComparableCheck,
+    (Concept::Comparable<InputImagePixelType>));
+  itkConceptMacro(OutputOStreamWritableCheck,
+    (Concept::OStreamWritable<OutputImagePixelType>));
+  /** End concept checking */
+#endif
+
+protected:
+  CudaZeroCrossingImageFilter()
+    {
+    m_ForegroundValue = NumericTraits<OutputImagePixelType>::One;
+    m_BackgroundValue = NumericTraits<OutputImagePixelType>::Zero;
+    m_CudaConf = CudaInterfaceType::New();
+    }
+  ~CudaZeroCrossingImageFilter(){}
+  void PrintSelf(std::ostream& os, Indent indent) const;
+  
+  CudaZeroCrossingImageFilter(const Self&) {}
+  OutputImagePixelType m_BackgroundValue;
+  OutputImagePixelType m_ForegroundValue;
+  
+  void GenerateData();
+  
+  typedef CudaInterface CudaInterfaceType;
+
+private:
+
+  typename CudaInterfaceType::Pointer m_CudaConf;
+
+};
+  
+} //end of namespace itk
+
+#ifndef ITK_MANUAL_INSTANTIATION
+#include "itkCudaZeroCrossingImageFilter.txx"
+#endif
+  
+#endif
diff -Naur CITK_lastest_devel_build/Code/BasicFilters/itkCudaZeroCrossingImageFilter.txx CITK_CudaCanny//Code/BasicFilters/itkCudaZeroCrossingImageFilter.txx
--- CITK_lastest_devel_build/Code/BasicFilters/itkCudaZeroCrossingImageFilter.txx	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/BasicFilters/itkCudaZeroCrossingImageFilter.txx	2011-08-07 19:17:50.349186000 -0300
@@ -0,0 +1,76 @@
+/*=========================================================================
+
+  Program:   Insight Segmentation & Registration Toolkit
+  Module:    $RCSfile: itkCudaZeroCrossingImageFilter.txx,v $
+  Language:  C++
+  Date:      $Date: 2010-11-04 17:31:03 $
+  Version:   $Revision: 0.1 $
+
+  Copyright (c) Insight Software Consortium. All rights reserved.
+  See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+     This software is distributed WITHOUT ANY WARRANTY; without even 
+     the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+     PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class CudaZeroCrossingImageFilter
+ * \brief This filter finds the closest pixel to the zero-crossings in a Image. 
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+#ifndef __itkCudaZeroCrossingImageFilter_txx
+#define __itkCudaZeroCrossingImageFilter_txx
+
+#include "itkCudaZeroCrossingImageFilter.h"
+
+#define THREADS_PER_BLOCK 256 
+
+namespace itk
+{
+
+template< class TInputImage, class TOutputImage >
+void
+CudaZeroCrossingImageFilter< TInputImage, TOutputImage >
+::GenerateData()
+{
+  // Set input, output and temporary pointers
+  typename TInputImage::ConstPointer input  = this->GetInput();
+  typename TOutputImage::Pointer output  = this->GetOutput();
+  typename TOutputImage::PixelType * tmp;
+  
+  m_CudaConf->SetBlockDim(THREADS_PER_BLOCK,1,1);
+  m_CudaConf->SetGridDim((this->GetInput()->GetPixelContainer()->Size()+THREADS_PER_BLOCK-1)/THREADS_PER_BLOCK,1,1);
+
+  // Allocate output image object
+  output->SetBufferedRegion(output->GetRequestedRegion());
+
+  // Get image size
+  typename OutputImageType::SizeType size = output->GetLargestPossibleRegion().GetSize();
+
+  // Call cudaSobel. Defined on CudaSobelEdgeDetection.cu
+  tmp = cudaZeroCrossing(m_CudaConf->GetGridDim(),m_CudaConf->GetBlockDim(),input->GetDevicePointer(), size[0], size[1]);
+
+  // Set image pointer to the output image
+  output->GetPixelContainer()->SetDevicePointer(tmp, size[0]*size[1], true);
+
+}
+
+template< class TInputImage, class TOutputImage >
+void
+CudaZeroCrossingImageFilter< TInputImage, TOutputImage >
+::PrintSelf( std::ostream& os, Indent indent ) const
+{
+  Superclass::PrintSelf(os,indent);
+
+  os << indent << "ForegroundValue: "
+     << static_cast<typename NumericTraits<OutputImagePixelType>::PrintType>(m_ForegroundValue)
+     << std::endl;
+  os << indent << "BackgroundValue: "
+     << static_cast<typename NumericTraits<OutputImagePixelType>::PrintType>(m_BackgroundValue)
+     << std::endl;
+}
+
+}//end of itk namespace
+
+#endif
diff -Naur CITK_lastest_devel_build/Code/Common/CMakeLists.txt CITK_CudaCanny//Code/Common/CMakeLists.txt
--- CITK_lastest_devel_build/Code/Common/CMakeLists.txt	2010-05-07 11:09:30.000000000 -0300
+++ CITK_CudaCanny//Code/Common/CMakeLists.txt	2011-08-07 19:04:20.085186000 -0300
@@ -42,7 +42,11 @@
   itkNumericTraitsRGBPixel.cxx
   itkNumericTraitsRGBAPixel.cxx
   itkNumericTraitsVectorPixel.cxx
+  itkNumericTraitsTensorPixel.cxx
+  itkNumericTraitsTensorPixel2.cxx
+  itkNumericTraitsDiffusionTensor3DPixel.cxx
   itkNumericTraitsFixedArrayPixel.cxx
+  itkNumericTraitsFixedArrayPixel2.cxx
   itkObject.cxx
   itkObjectFactoryBase.cxx
   itkOneWayEquivalencyTable.cxx
@@ -166,17 +170,26 @@
     ARCHIVE DESTINATION ${ITK_INSTALL_LIB_DIR_CM24} COMPONENT Development)
 ENDIF(NOT ITK_INSTALL_NO_LIBRARIES)
 
+IF(ITK_USE_FLAT_DIRECTORY_INSTALL)
+  SET( ITK_INSTALL_INCLUDE_DIR_CM24_COMMON ${ITK_INSTALL_INCLUDE_DIR_CM24})
+  SET( ITK_INSTALL_INCLUDE_DIR_CM24_COMMON_TEMPLATES ${ITK_INSTALL_INCLUDE_DIR_CM24})
+ELSE(ITK_USE_FLAT_DIRECTORY_INSTALL)
+  SET( ITK_INSTALL_INCLUDE_DIR_CM24_COMMON ${ITK_INSTALL_INCLUDE_DIR_CM24}/Common)
+  SET( ITK_INSTALL_INCLUDE_DIR_CM24_COMMON_TEMPLATES ${ITK_INSTALL_INCLUDE_DIR_CM24}/Common/Templates)
+ENDIF(ITK_USE_FLAT_DIRECTORY_INSTALL)
+
+
 IF(NOT ITK_INSTALL_NO_DEVELOPMENT)
   FILE(GLOB __files1 "${CMAKE_CURRENT_SOURCE_DIR}/*.h")
   FILE(GLOB __files2 "${CMAKE_CURRENT_SOURCE_DIR}/*.txx")
   INSTALL(FILES ${__files1} ${__files2}
-    DESTINATION ${ITK_INSTALL_INCLUDE_DIR_CM24}/Common
+    DESTINATION ${ITK_INSTALL_INCLUDE_DIR_CM24_COMMON}
     COMPONENT Development)
 
   IF(ITK_EXPLICIT_INSTANTIATION)
     FILE(GLOB __files3 "${CMAKE_CURRENT_BINARY_DIR}/Templates/*.h")
     INSTALL(FILES ${__files3}
-      DESTINATION ${ITK_INSTALL_INCLUDE_DIR_CM24}/Common/Templates
+      DESTINATION ${ITK_INSTALL_INCLUDE_DIR_CM24_COMMON_TEMPLATES}
       COMPONENT Development)
   ENDIF(ITK_EXPLICIT_INSTANTIATION)
 ENDIF(NOT ITK_INSTALL_NO_DEVELOPMENT)
diff -Naur CITK_lastest_devel_build/Code/Common/itkCudaImportImageContainer.h CITK_CudaCanny//Code/Common/itkCudaImportImageContainer.h
--- CITK_lastest_devel_build/Code/Common/itkCudaImportImageContainer.h	2010-05-07 11:09:30.000000000 -0300
+++ CITK_CudaCanny//Code/Common/itkCudaImportImageContainer.h	2011-08-07 19:04:20.085186000 -0300
@@ -74,7 +74,7 @@
          TElement *GetImportPointer() const
          {
             if (ImageLocation==GPU) { CopyToCPU(); ImageLocation=BOTH; }
-            return m_DevicePointer;
+            return m_ImportPointer;
          };
 
          TElement *GetDevicePointer() 
diff -Naur CITK_lastest_devel_build/Code/Common/itkCudaInterface.h CITK_CudaCanny//Code/Common/itkCudaInterface.h
--- CITK_lastest_devel_build/Code/Common/itkCudaInterface.h	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/Common/itkCudaInterface.h	2011-08-07 19:18:17.369186000 -0300
@@ -0,0 +1,108 @@
+/*=========================================================================
+
+Program:   Insight Segmentation & Registration Toolkit
+Module:    $RCSfile: itkCudaInterface.h,v $
+Language:  C++
+Date:      $Date: 2010-11-19 12:24:09 $
+Version:   $Revision: 0.1 $
+
+Copyright (c) Insight Software Consortium. All rights reserved.
+See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+This software is distributed WITHOUT ANY WARRANTY; without even 
+the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class itkCudaInterface
+ * \brief Class to store CUDA settings
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#ifndef __itkCudaInterface_h
+#define __itkCudaInterface_h
+
+#include <cuda.h>
+#include <cutil_inline.h>
+#include <utility>
+
+namespace itk
+{
+
+    class ITK_EXPORT CudaInterface :
+      public Object
+   {
+      public:
+         /** Standard class typedefs. */
+         typedef CudaInterface    Self;
+         typedef Object                    Superclass;
+         typedef SmartPointer<Self>        Pointer;
+         typedef SmartPointer<const Self>  ConstPointer;
+
+         /** Run-time type information (and related methods). */
+         itkTypeMacro(CudaInterface,Object);
+
+         /** Method for creation through the object factory. */
+         itkNewMacro(Self);
+
+         /** Get the cuda block and grid dimensions */
+         dim3 GetBlockDim() 
+         { 
+            return m_BlockDim;
+         };
+
+         dim3 GetGridDim()
+         {
+            return m_GridDim;
+         };
+
+         /** Set the cuda block and grid dimensions */
+         void SetBlockDim(uint x, uint y, uint z)
+         {
+           m_BlockDim.x = x;
+           m_BlockDim.y = y;
+           m_BlockDim.z = z;
+         };
+
+         void SetGridDim(uint x, uint y, uint z)
+         {
+           m_GridDim.x = x;
+           m_GridDim.y = y;
+           m_GridDim.z = 1;
+         };
+
+
+      protected:
+         CudaInterface();
+         virtual ~CudaInterface(){};
+
+         /** PrintSelf routine. Normally this is a protected internal method. It is
+          * made public here so that Image can call this method.  Users should not
+          * call this method but should call Print() instead. */
+         void PrintSelf(std::ostream& os, Indent indent) const;
+
+      private:
+         CudaInterface(const Self&); //purposely not implemented
+
+         dim3 m_BlockDim;
+         dim3 m_GridDim;
+   };
+
+} // end namespace itk
+
+// Define instantiation macro for this template.
+#define ITK_TEMPLATE_CudaInterface(_, EXPORT, x, y) namespace itk { \
+   _(2(class EXPORT CudaInterface< ITK_TEMPLATE_2 x >)) \
+   namespace Templates { typedef CudaInterface<ITK_TEMPLATE_2 x > CudaInterface##y; } \
+}
+
+#if ITK_TEMPLATE_EXPLICIT
+# include "Templates/itkCudaInterface+-.h"
+#endif
+
+#if ITK_TEMPLATE_TXX
+# include "itkCudaInterface.txx"
+#endif
+
+#endif
diff -Naur CITK_lastest_devel_build/Code/Common/itkCudaInterface.txx CITK_CudaCanny//Code/Common/itkCudaInterface.txx
--- CITK_lastest_devel_build/Code/Common/itkCudaInterface.txx	1969-12-31 21:00:00.000000000 -0300
+++ CITK_CudaCanny//Code/Common/itkCudaInterface.txx	2011-08-07 19:18:17.437186000 -0300
@@ -0,0 +1,65 @@
+/*=========================================================================
+
+Program:   Insight Segmentation & Registration Toolkit
+Module:    $RCSfile: itkCudaInterface.txx,v $
+Language:  C++
+Date:      $Date: 2010-11-19 19:10:47 $
+Version:   $Revision: 0.1 $
+
+Copyright (c) Insight Software Consortium. All rights reserved.
+See ITKCopyright.txt or http://www.itk.org/HTML/Copyright.htm for details.
+
+Portions of this code are covered under the VTK copyright.
+See VTKCopyright.txt or http://www.kitware.com/VTKCopyright.htm for details.
+
+This software is distributed WITHOUT ANY WARRANTY; without even 
+the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR 
+PURPOSE.  See the above copyright notices for more information.
+
+=========================================================================*/
+/** \class itkCudaInterface
+ * \brief Class to store CUDA settings
+ * \author Luis Lourenco - Vision, Robotics, and Image Group (VRI)
+ * http://www.inf.ufpr.br/vri/alumni/2011-LuisLourenco/
+**/
+
+#ifndef __itkCudaInterface_txx
+#define __itkCudaInterface_txx
+
+#include "itkCudaInterface.h"
+#include <cstring>
+#include <stdlib.h>
+#include <string.h>
+
+namespace itk
+{
+
+      CudaInterface
+      ::CudaInterface()
+      {
+        m_BlockDim.x = 256;
+        m_BlockDim.y = 1;
+        m_BlockDim.z = 1;
+        m_GridDim.x = 1;
+        m_GridDim.y = 1;
+        m_GridDim.z = 1;
+      }
+
+
+
+      void 
+      CudaInterface
+      ::PrintSelf(std::ostream& os, Indent indent) const
+      {
+         os << indent << "BlockDim.x: " << m_BlockDim.x << std::endl;
+         os << indent << "BlockDim.y: " << m_BlockDim.y << std::endl;
+         os << indent << "BlockDim.z: " << m_BlockDim.z << std::endl;
+         os << indent << "GridDim.x: " << m_GridDim.x << std::endl;
+         os << indent << "GridDim.y: " << m_GridDim.y << std::endl;
+         os << indent << "GridDim.z: " << m_GridDim.z << std::endl;
+      }
+
+
+} // end namespace itk
+
+#endif
