From e3452670bffd63d93f7abf2a354694f37be2fe97 Mon Sep 17 00:00:00 2001 From: Axel Garcia Date: Fri, 21 Nov 2025 15:15:07 +0100 Subject: [PATCH 1/2] ENH: Add CudaSquareImageFilter types and wrap it --- include/itkCudaSquareImage.hcu | 4 ++ include/itkCudaSquareImageFilter.hxx | 25 +++++--- src/itkCudaSquareImage.cu | 61 ++++++++++++++++++- wrapping/CMakeLists.txt | 15 +++++ wrapping/itkCudaImageToImageFilter.wrap | 43 +++++++++++++ wrapping/itkCudaInPlaceImageFilter.wrap | 23 +++++++ wrapping/itkCudaSquareImageFilter.wrap | 13 ++++ wrapping/itkInPlaceImageFilterCudaCommon.wrap | 22 +++++++ 8 files changed, 198 insertions(+), 8 deletions(-) create mode 100644 wrapping/itkCudaImageToImageFilter.wrap create mode 100644 wrapping/itkCudaInPlaceImageFilter.wrap create mode 100644 wrapping/itkCudaSquareImageFilter.wrap create mode 100644 wrapping/itkInPlaceImageFilterCudaCommon.wrap diff --git a/include/itkCudaSquareImage.hcu b/include/itkCudaSquareImage.hcu index 0af2660..74baffe 100644 --- a/include/itkCudaSquareImage.hcu +++ b/include/itkCudaSquareImage.hcu @@ -25,6 +25,10 @@ template void CudaSquareImage3D(int imSize[3], PixelType * in, PixelType * out); +template +void +CudaSquareImage2D(int imSize[2], PixelType * in, PixelType * out); + } // end namespace itk #endif diff --git a/include/itkCudaSquareImageFilter.hxx b/include/itkCudaSquareImageFilter.hxx index 4ec5085..5913396 100644 --- a/include/itkCudaSquareImageFilter.hxx +++ b/include/itkCudaSquareImageFilter.hxx @@ -26,18 +26,29 @@ template void CudaSquareImageFilter::GPUGenerateData() { - int size[3] = { 1, 1, 1 }; - for (unsigned int i = 0; i < ImageDimension; i++) - { - size[i] = this->GetInput()->GetBufferedRegion().GetSize()[i]; - } - typename ImageType::PixelType * pin0 = (typename ImageType::PixelType *)(this->GetInput(0)->GetCudaDataManager()->GetGPUBufferPointer()); typename ImageType::PixelType * pout = (typename ImageType::PixelType *)(this->GetOutput()->GetCudaDataManager()->GetGPUBufferPointer()); - CudaSquareImage3D(size, pin0, pout); + if constexpr (ImageDimension == 2) + { + int size[2]; + for (unsigned int i = 0; i < 2; ++i) + { + size[i] = this->GetInput()->GetBufferedRegion().GetSize()[i]; + } + CudaSquareImage2D(size, pin0, pout); + } + if constexpr (ImageDimension == 3) + { + int size[3] = { 1, 1, 1 }; + for (unsigned int i = 0; i < ImageDimension; ++i) + { + size[i] = this->GetInput()->GetBufferedRegion().GetSize()[i]; + } + CudaSquareImage3D(size, pin0, pout); + } } } // end namespace itk diff --git a/src/itkCudaSquareImage.cu b/src/itkCudaSquareImage.cu index 3f9ece1..6efed25 100644 --- a/src/itkCudaSquareImage.cu +++ b/src/itkCudaSquareImage.cu @@ -21,6 +21,42 @@ namespace itk { +template +__global__ void +CudaSquareImage2D_kernel(int2 imSize, PixelType * in, PixelType * out) +{ + unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; + unsigned int j = blockIdx.y * blockDim.y + threadIdx.y; + + if (i >= imSize.x || j >= imSize.y) + { + return; + } + unsigned int gidx = i + j * imSize.x; + + out[gidx] = in[gidx] * in[gidx]; +} + +template +void +CudaSquareImage2D(int imSize[2], PixelType * in, PixelType * out) +{ + // Thread Block Dimensions + constexpr int tBlock_x = 16; + constexpr int tBlock_y = 16; + + unsigned int blocksInX = (imSize[0] - 1) / tBlock_x + 1; + unsigned int blocksInY = (imSize[1] - 1) / tBlock_y + 1; + + // Compute block and grid sizes + dim3 dimGrid = dim3(blocksInX, blocksInY); + dim3 dimBlock = dim3(tBlock_x, tBlock_y); + + int2 imageSize = make_int2(imSize[0], imSize[1]); + + CudaSquareImage2D_kernel<<>>(imageSize, in, out); +} + template __global__ void CudaSquareImage3D_kernel(int3 imSize, PixelType * in, PixelType * out) @@ -64,5 +100,28 @@ template void CudaCommon_EXPORT CudaSquareImage3D(int imSize[3], float * in, float * out); template void CudaCommon_EXPORT CudaSquareImage3D(int imSize[3], double * in, double * out); - +template void CudaCommon_EXPORT +CudaSquareImage3D(int imSize[3], int * in, int * out); +template void CudaCommon_EXPORT +CudaSquareImage3D(int imSize[3], unsigned int * in, unsigned int * out); +template void CudaCommon_EXPORT +CudaSquareImage3D(int imSize[3], short * in, short * out); +template void CudaCommon_EXPORT +CudaSquareImage3D(int imSize[3], unsigned short * in, unsigned short * out); +template void CudaCommon_EXPORT +CudaSquareImage3D(int imSize[3], unsigned char * in, unsigned char * out); +template void CudaCommon_EXPORT +CudaSquareImage2D(int imSize[2], float * in, float * out); +template void CudaCommon_EXPORT +CudaSquareImage2D(int imSize[2], double * in, double * out); +template void CudaCommon_EXPORT +CudaSquareImage2D(int imSize[2], int * in, int * out); +template void CudaCommon_EXPORT +CudaSquareImage2D(int imSize[2], unsigned int * in, unsigned int * out); +template void CudaCommon_EXPORT +CudaSquareImage2D(int imSize[2], short * in, short * out); +template void CudaCommon_EXPORT +CudaSquareImage2D(int imSize[2], unsigned short * in, unsigned short * out); +template void CudaCommon_EXPORT +CudaSquareImage2D(int imSize[2], unsigned char * in, unsigned char * out); } // end namespace itk diff --git a/wrapping/CMakeLists.txt b/wrapping/CMakeLists.txt index a76d2d0..1e8f59d 100644 --- a/wrapping/CMakeLists.txt +++ b/wrapping/CMakeLists.txt @@ -1,4 +1,19 @@ itk_wrap_module(CudaCommon) + +set( + WRAPPER_SUBMODULE_ORDER + itkCudaDataManager + itkCudaImage + itkCudaImageDataManager + itkCudaImageFromImageFilter + itkCudaImageToImageFilter + itkCudaInPlaceImageFilter + itkCudaSquareImageFilter + itkImageSourceCudaCommon + itkImageToImageFilterCudaCommon + itkInPlaceImageFilterCudaCommon +) + itk_auto_load_submodules() itk_end_wrap_module() diff --git a/wrapping/itkCudaImageToImageFilter.wrap b/wrapping/itkCudaImageToImageFilter.wrap new file mode 100644 index 0000000..8adf5d6 --- /dev/null +++ b/wrapping/itkCudaImageToImageFilter.wrap @@ -0,0 +1,43 @@ +itk_wrap_include(itkCudaImage.h) +itk_wrap_include(itkCudaImageToImageFilter.h) +itk_wrap_include(itkInPlaceImageFilter.h) + +itk_wrap_class("itk::CudaImageToImageFilter" POINTER) + + # Match CudaImage scalar instantiations + UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(t ${types}) + itk_wrap_template("CI${ITKM_${t}}${d}CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>") + endforeach() + endforeach() + + # Match CudaImage vector/covariant instantiations + UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") + foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(vt ${vector_types}) + itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + endforeach() + endforeach() + endforeach() + + # Instantiate with InPlaceImageFilter as parent for in-place CUDA filters + UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(t ${types}) + itk_wrap_template("CI${ITKM_${t}}${d}CI${ITKM_${t}}${d}IPF" "itk::CudaImage<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>, itk::InPlaceImageFilter, itk::CudaImage<${ITKT_${t}}, ${d}> >") + endforeach() + endforeach() + + UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") + foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(vt ${vector_types}) + itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}IPF" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::InPlaceImageFilter, itk::CudaImage<${ITKT_${vt}${c}}, ${d}> >") + endforeach() + endforeach() + endforeach() + + +itk_end_wrap_class() diff --git a/wrapping/itkCudaInPlaceImageFilter.wrap b/wrapping/itkCudaInPlaceImageFilter.wrap new file mode 100644 index 0000000..793f679 --- /dev/null +++ b/wrapping/itkCudaInPlaceImageFilter.wrap @@ -0,0 +1,23 @@ +itk_wrap_include(itkCudaImage.h) + +itk_wrap_class("itk::CudaInPlaceImageFilter" POINTER) + + UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(t ${types}) + itk_wrap_template("CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>") + itk_wrap_template("CI${ITKM_${t}}${d}CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>") + endforeach() + endforeach() + + UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") + foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(vt ${vector_types}) + itk_wrap_template("CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + endforeach() + endforeach() + endforeach() + +itk_end_wrap_class() diff --git a/wrapping/itkCudaSquareImageFilter.wrap b/wrapping/itkCudaSquareImageFilter.wrap new file mode 100644 index 0000000..3874579 --- /dev/null +++ b/wrapping/itkCudaSquareImageFilter.wrap @@ -0,0 +1,13 @@ +itk_wrap_include(itkCudaImage.h) + +itk_wrap_class("itk::CudaSquareImageFilter" POINTER) + + # Restrict wrapping to the scalar types we explicitly instantiate in itkCudaSquareImage.cu + UNIQUE(types "F;D;UC") + foreach(t ${types}) + foreach(d 2 3) + itk_wrap_template("CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}},${d}>") + endforeach() + endforeach() + +itk_end_wrap_class() diff --git a/wrapping/itkInPlaceImageFilterCudaCommon.wrap b/wrapping/itkInPlaceImageFilterCudaCommon.wrap new file mode 100644 index 0000000..27e266f --- /dev/null +++ b/wrapping/itkInPlaceImageFilterCudaCommon.wrap @@ -0,0 +1,22 @@ +itk_wrap_include(itkCudaImage.h) +itk_wrap_include(itkInPlaceImageFilter.h) + +itk_wrap_class("itk::InPlaceImageFilter" POINTER) + + UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(t ${types}) + itk_wrap_template("IPCI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>") + endforeach() + endforeach() + + UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") + foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(vt ${vector_types}) + itk_wrap_template("IPCI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + endforeach() + endforeach() + endforeach() + +itk_end_wrap_class() From fb45528f23d617ccdffb98ed196a38cfc4f4a768 Mon Sep 17 00:00:00 2001 From: Axel Garcia Date: Tue, 27 Jan 2026 10:09:11 +0100 Subject: [PATCH 2/2] ENH: Adapt cuda filters wrapping to follow itk::GPUImage --- wrapping/CMakeLists.txt | 38 +++++++++++++++++-- wrapping/itkCudaImage.wrap | 2 +- wrapping/itkCudaImageDataManager.wrap | 2 +- wrapping/itkCudaImageFromImageFilter.wrap | 2 +- wrapping/itkCudaImageToImageFilter.wrap | 38 +++---------------- wrapping/itkCudaInPlaceImageFilter.wrap | 24 ++++-------- wrapping/itkImageSourceCudaCommon.wrap | 4 +- wrapping/itkImageToImageFilterCudaCommon.wrap | 31 +++++++++------ wrapping/itkInPlaceImageFilterCudaCommon.wrap | 29 +++++++++----- 9 files changed, 94 insertions(+), 76 deletions(-) diff --git a/wrapping/CMakeLists.txt b/wrapping/CMakeLists.txt index 1e8f59d..44ab6b6 100644 --- a/wrapping/CMakeLists.txt +++ b/wrapping/CMakeLists.txt @@ -1,3 +1,33 @@ +macro(cuda_wrap_image_filter_combinations from_types to_types prefix) + set(parent ${ARGN}) + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(from ${from_types}) + foreach(to ${to_types}) + if(parent) + itk_wrap_template("${prefix}CI${ITKM_${from}}${d}CI${ITKM_${to}}${d}" "itk::CudaImage<${ITKT_${from}}, ${d}>, itk::CudaImage<${ITKT_${to}}, ${d}>, ${parent}, itk::CudaImage<${ITKT_${to}}, ${d}> >") + else() + itk_wrap_template("${prefix}CI${ITKM_${from}}${d}CI${ITKM_${to}}${d}" "itk::CudaImage<${ITKT_${from}}, ${d}>, itk::CudaImage<${ITKT_${to}}, ${d}>") + endif() + endforeach() + endforeach() + endforeach() +endmacro() + +macro(cuda_wrap_vector_combinations vector_types prefix) + set(parent ${ARGN}) + foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) + foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(vt ${vector_types}) + if(parent) + itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}${prefix}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, ${parent}, itk::CudaImage<${ITKT_${vt}${c}}, ${d}> >") + else() + itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}${prefix}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + endif() + endforeach() + endforeach() + endforeach() +endmacro() + itk_wrap_module(CudaCommon) set( @@ -5,13 +35,13 @@ set( itkCudaDataManager itkCudaImage itkCudaImageDataManager - itkCudaImageFromImageFilter + itkInPlaceImageFilterCudaCommon + itkImageSourceCudaCommon + itkImageToImageFilterCudaCommon itkCudaImageToImageFilter itkCudaInPlaceImageFilter + itkCudaImageFromImageFilter itkCudaSquareImageFilter - itkImageSourceCudaCommon - itkImageToImageFilterCudaCommon - itkInPlaceImageFilterCudaCommon ) itk_auto_load_submodules() diff --git a/wrapping/itkCudaImage.wrap b/wrapping/itkCudaImage.wrap index e24e37d..cc62179 100644 --- a/wrapping/itkCudaImage.wrap +++ b/wrapping/itkCudaImage.wrap @@ -1,7 +1,7 @@ configure_file("${CMAKE_CURRENT_SOURCE_DIR}/CudaImage.i.init" "${CMAKE_CURRENT_BINARY_DIR}/CudaImage.i" @ONLY) itk_wrap_class("itk::CudaImage" POINTER_WITH_CONST_POINTER) - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + UNIQUE(types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") UNIQUE(vector_universe "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") foreach(t ${types}) set(PixelType ${t}) diff --git a/wrapping/itkCudaImageDataManager.wrap b/wrapping/itkCudaImageDataManager.wrap index 7d61f28..aee99cb 100644 --- a/wrapping/itkCudaImageDataManager.wrap +++ b/wrapping/itkCudaImageDataManager.wrap @@ -2,7 +2,7 @@ itk_wrap_include(itkCudaImage.h) itk_wrap_class("itk::CudaImageDataManager" POINTER) - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + UNIQUE(types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") foreach(d ${ITK_WRAP_IMAGE_DIMS}) foreach(t ${types}) itk_wrap_template("CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>") diff --git a/wrapping/itkCudaImageFromImageFilter.wrap b/wrapping/itkCudaImageFromImageFilter.wrap index 9cff53c..9b4617e 100644 --- a/wrapping/itkCudaImageFromImageFilter.wrap +++ b/wrapping/itkCudaImageFromImageFilter.wrap @@ -2,7 +2,7 @@ itk_wrap_include(itkImage.h) itk_wrap_class("itk::CudaImageFromImageFilter" POINTER) - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + UNIQUE(types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") foreach(d ${ITK_WRAP_IMAGE_DIMS}) foreach(t ${types}) itk_wrap_template("${t}${d}" "itk::Image<${ITKT_${t}}, ${d}>") diff --git a/wrapping/itkCudaImageToImageFilter.wrap b/wrapping/itkCudaImageToImageFilter.wrap index 8adf5d6..8dc88d5 100644 --- a/wrapping/itkCudaImageToImageFilter.wrap +++ b/wrapping/itkCudaImageToImageFilter.wrap @@ -4,40 +4,14 @@ itk_wrap_include(itkInPlaceImageFilter.h) itk_wrap_class("itk::CudaImageToImageFilter" POINTER) - # Match CudaImage scalar instantiations - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") - foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(t ${types}) - itk_wrap_template("CI${ITKM_${t}}${d}CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>") - endforeach() - endforeach() - - # Match CudaImage vector/covariant instantiations + unique(from_types "D;UC;UL;${WRAP_ITK_SCALAR}") + unique(to_types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") - foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) - foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(vt ${vector_types}) - itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") - endforeach() - endforeach() - endforeach() - - # Instantiate with InPlaceImageFilter as parent for in-place CUDA filters - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") - foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(t ${types}) - itk_wrap_template("CI${ITKM_${t}}${d}CI${ITKM_${t}}${d}IPF" "itk::CudaImage<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>, itk::InPlaceImageFilter, itk::CudaImage<${ITKT_${t}}, ${d}> >") - endforeach() - endforeach() - UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") - foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) - foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(vt ${vector_types}) - itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}IPF" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::InPlaceImageFilter, itk::CudaImage<${ITKT_${vt}${c}}, ${d}> >") - endforeach() - endforeach() - endforeach() + cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "" "") + cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "IP" "itk::InPlaceImageFilter") + cuda_wrap_vector_combinations("${vector_types}" "" "") + cuda_wrap_vector_combinations("${vector_types}" "IP" "itk::InPlaceImageFilter") itk_end_wrap_class() diff --git a/wrapping/itkCudaInPlaceImageFilter.wrap b/wrapping/itkCudaInPlaceImageFilter.wrap index 793f679..a584492 100644 --- a/wrapping/itkCudaInPlaceImageFilter.wrap +++ b/wrapping/itkCudaInPlaceImageFilter.wrap @@ -2,22 +2,14 @@ itk_wrap_include(itkCudaImage.h) itk_wrap_class("itk::CudaInPlaceImageFilter" POINTER) - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") - foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(t ${types}) - itk_wrap_template("CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>") - itk_wrap_template("CI${ITKM_${t}}${d}CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>") - endforeach() - endforeach() - + unique(from_types "D;UC;UL;${WRAP_ITK_SCALAR}") + unique(to_types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") - foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) - foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(vt ${vector_types}) - itk_wrap_template("CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") - itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") - endforeach() - endforeach() - endforeach() + + cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "IP") + cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "IP" "itk::InPlaceImageFilter") + + cuda_wrap_vector_combinations("${vector_types}" "IP") + cuda_wrap_vector_combinations("${vector_types}" "IP" "itk::InPlaceImageFilter") itk_end_wrap_class() diff --git a/wrapping/itkImageSourceCudaCommon.wrap b/wrapping/itkImageSourceCudaCommon.wrap index 47d0aa1..bca3d2c 100644 --- a/wrapping/itkImageSourceCudaCommon.wrap +++ b/wrapping/itkImageSourceCudaCommon.wrap @@ -2,10 +2,11 @@ itk_wrap_include(itkCudaImage.h) itk_wrap_class("itk::ImageSource" POINTER) - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + UNIQUE(types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") foreach(d ${ITK_WRAP_IMAGE_DIMS}) foreach(t ${types}) itk_wrap_template("CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>") + itk_wrap_template("I${ITKM_${t}}${d}" "itk::Image<${ITKT_${t}}, ${d}>") endforeach() endforeach() @@ -14,6 +15,7 @@ itk_wrap_class("itk::ImageSource" POINTER) foreach(d ${ITK_WRAP_IMAGE_DIMS}) foreach(vt ${vector_types}) itk_wrap_template("CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + itk_wrap_template("I${ITKM_${vt}${c}}${d}" "itk::Image<${ITKT_${vt}${c}}, ${d}>") endforeach() endforeach() endforeach() diff --git a/wrapping/itkImageToImageFilterCudaCommon.wrap b/wrapping/itkImageToImageFilterCudaCommon.wrap index a31c86f..cef7f7c 100644 --- a/wrapping/itkImageToImageFilterCudaCommon.wrap +++ b/wrapping/itkImageToImageFilterCudaCommon.wrap @@ -3,20 +3,29 @@ itk_wrap_include(itkImage.h) itk_wrap_class("itk::ImageToImageFilter" POINTER) - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + unique(from_types "D;UC;UL;${WRAP_ITK_SCALAR}") + unique(to_types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") + foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(t ${types}) - itk_wrap_template("CI${ITKM_${t}}${d}CI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>") - itk_wrap_template("I${ITKM_${t}}${d}CI${ITKM_${t}}${d}" "itk::Image<${ITKT_${t}}, ${d}>, itk::CudaImage<${ITKT_${t}}, ${d}>") + foreach(from ${from_types}) + foreach(to ${to_types}) + itk_wrap_template("I${ITKM_${from}}${d}CI${ITKM_${to}}${d}" + "itk::Image< ${ITKT_${from}}, ${d} >, itk::CudaImage< ${ITKT_${to}}, ${d} >") + itk_wrap_template("CI${ITKM_${from}}${d}I${ITKM_${to}}${d}" + "itk::CudaImage< ${ITKT_${from}}, ${d} >, itk::Image< ${ITKT_${to}}, ${d} >") + itk_wrap_template("CI${ITKM_${from}}${d}CI${ITKM_${to}}${d}" + "itk::CudaImage< ${ITKT_${from}}, ${d} >, itk::CudaImage< ${ITKT_${to}}, ${d} >") + endforeach() endforeach() - endforeach() - - UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") - foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) - foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) foreach(vt ${vector_types}) - itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") - itk_wrap_template("I${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" "itk::Image<${ITKT_${vt}${c}}, ${d}>, itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + itk_wrap_template("I${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" + "itk::Image< ${ITKT_${vt}${c}}, ${d} >, itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >") + itk_wrap_template("CI${ITKM_${vt}${c}}${d}I${ITKM_${vt}${c}}${d}" + "itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >, itk::Image< ${ITKT_${vt}${c}}, ${d} >") + itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" + "itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >, itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >") endforeach() endforeach() endforeach() diff --git a/wrapping/itkInPlaceImageFilterCudaCommon.wrap b/wrapping/itkInPlaceImageFilterCudaCommon.wrap index 27e266f..35bdd2f 100644 --- a/wrapping/itkInPlaceImageFilterCudaCommon.wrap +++ b/wrapping/itkInPlaceImageFilterCudaCommon.wrap @@ -3,18 +3,29 @@ itk_wrap_include(itkInPlaceImageFilter.h) itk_wrap_class("itk::InPlaceImageFilter" POINTER) - UNIQUE(types "UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + unique(from_types "D;UC;UL;${WRAP_ITK_SCALAR}") + unique(to_types "D;UC;UL;${ITKM_IT};${WRAP_ITK_SCALAR}") + unique(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") + foreach(d ${ITK_WRAP_IMAGE_DIMS}) - foreach(t ${types}) - itk_wrap_template("IPCI${ITKM_${t}}${d}" "itk::CudaImage<${ITKT_${t}}, ${d}>") + foreach(from ${from_types}) + foreach(to ${to_types}) + itk_wrap_template("I${ITKM_${from}}${d}CI${ITKM_${to}}${d}" + "itk::Image< ${ITKT_${from}}, ${d} >, itk::CudaImage< ${ITKT_${to}}, ${d} >") + itk_wrap_template("CI${ITKM_${from}}${d}I${ITKM_${to}}${d}" + "itk::CudaImage< ${ITKT_${from}}, ${d} >, itk::Image< ${ITKT_${to}}, ${d} >") + itk_wrap_template("CI${ITKM_${from}}${d}CI${ITKM_${to}}${d}" + "itk::CudaImage< ${ITKT_${from}}, ${d} >, itk::CudaImage< ${ITKT_${to}}, ${d} >") + endforeach() endforeach() - endforeach() - - UNIQUE(vector_types "${WRAP_ITK_VECTOR_REAL};${WRAP_ITK_COV_VECTOR_REAL}") - foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) - foreach(d ${ITK_WRAP_IMAGE_DIMS}) + foreach(c ${ITK_WRAP_VECTOR_COMPONENTS}) foreach(vt ${vector_types}) - itk_wrap_template("IPCI${ITKM_${vt}${c}}${d}" "itk::CudaImage<${ITKT_${vt}${c}}, ${d}>") + itk_wrap_template("I${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" + "itk::Image< ${ITKT_${vt}${c}}, ${d} >, itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >") + itk_wrap_template("CI${ITKM_${vt}${c}}${d}I${ITKM_${vt}${c}}${d}" + "itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >, itk::Image< ${ITKT_${vt}${c}}, ${d} >") + itk_wrap_template("CI${ITKM_${vt}${c}}${d}CI${ITKM_${vt}${c}}${d}" + "itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >, itk::CudaImage< ${ITKT_${vt}${c}}, ${d} >") endforeach() endforeach() endforeach()