Skip to content

CudaSquareImageFilter #89

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
axel-grc wants to merge 2 commits into
base: main
Choose a base branch
from
Open

CudaSquareImageFilter #89

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
e345267
ENH: Add CudaSquareImageFilter types and wrap it
axel-grc Nov 21, 2025
fb45528
ENH: Adapt cuda filters wrapping to follow itk::GPUImage
axel-grc Jan 27, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions include/itkCudaSquareImage.hcu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,10 @@ template <class PixelType>
void
CudaSquareImage3D(int imSize[3], PixelType * in, PixelType * out);

template <class PixelType>
void
CudaSquareImage2D(int imSize[2], PixelType * in, PixelType * out);

} // end namespace itk

#endif
25 changes: 18 additions & 7 deletions include/itkCudaSquareImageFilter.hxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,18 +26,29 @@ template <class ImageType>
void
CudaSquareImageFilter<ImageType>::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<typename ImageType::PixelType>(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<typename ImageType::PixelType>(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<typename ImageType::PixelType>(size, pin0, pout);
}
}

} // end namespace itk
Expand Down
61 changes: 60 additions & 1 deletion src/itkCudaSquareImage.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,42 @@
namespace itk
{

template <class PixelType>
__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 <class PixelType>
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<PixelType><<<dimGrid, dimBlock>>>(imageSize, in, out);
}

template <class PixelType>
__global__ void
CudaSquareImage3D_kernel(int3 imSize, PixelType * in, PixelType * out)
Expand Down Expand Up @@ -64,5 +100,28 @@ template void CudaCommon_EXPORT
CudaSquareImage3D<float>(int imSize[3], float * in, float * out);
template void CudaCommon_EXPORT
CudaSquareImage3D<double>(int imSize[3], double * in, double * out);

template void CudaCommon_EXPORT
CudaSquareImage3D<int>(int imSize[3], int * in, int * out);
template void CudaCommon_EXPORT
CudaSquareImage3D<unsigned int>(int imSize[3], unsigned int * in, unsigned int * out);
template void CudaCommon_EXPORT
CudaSquareImage3D<short>(int imSize[3], short * in, short * out);
template void CudaCommon_EXPORT
CudaSquareImage3D<unsigned short>(int imSize[3], unsigned short * in, unsigned short * out);
template void CudaCommon_EXPORT
CudaSquareImage3D<unsigned char>(int imSize[3], unsigned char * in, unsigned char * out);
template void CudaCommon_EXPORT
CudaSquareImage2D<float>(int imSize[2], float * in, float * out);
template void CudaCommon_EXPORT
CudaSquareImage2D<double>(int imSize[2], double * in, double * out);
template void CudaCommon_EXPORT
CudaSquareImage2D<int>(int imSize[2], int * in, int * out);
template void CudaCommon_EXPORT
CudaSquareImage2D<unsigned int>(int imSize[2], unsigned int * in, unsigned int * out);
template void CudaCommon_EXPORT
CudaSquareImage2D<short>(int imSize[2], short * in, short * out);
template void CudaCommon_EXPORT
CudaSquareImage2D<unsigned short>(int imSize[2], unsigned short * in, unsigned short * out);
template void CudaCommon_EXPORT
CudaSquareImage2D<unsigned char>(int imSize[2], unsigned char * in, unsigned char * out);
} // end namespace itk
45 changes: 45 additions & 0 deletions wrapping/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,49 @@
macro(cuda_wrap_image_filter_combinations from_types to_types prefix)
set(parent ${ARGN})
foreach(d ${ITK_WRAP_IMAGE_DIMS})
foreach(from ${from_types})
Copy link
Collaborator

@SimonRit SimonRit Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we really combinations from different types as input and output?

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_${from}}, ${d}>, 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}>, 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(
WRAPPER_SUBMODULE_ORDER
itkCudaDataManager
itkCudaImage
itkCudaImageDataManager
itkInPlaceImageFilterCudaCommon
itkImageSourceCudaCommon
itkImageToImageFilterCudaCommon
itkCudaImageToImageFilter
itkCudaInPlaceImageFilter
itkCudaImageFromImageFilter
itkCudaSquareImageFilter
)

itk_auto_load_submodules()
itk_end_wrap_module()

Expand Down
2 changes: 1 addition & 1 deletion wrapping/itkCudaImage.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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})
Expand Down
2 changes: 1 addition & 1 deletion wrapping/itkCudaImageDataManager.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}>")
Expand Down
2 changes: 1 addition & 1 deletion wrapping/itkCudaImageFromImageFilter.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}>")
Expand Down
17 changes: 17 additions & 0 deletions wrapping/itkCudaImageToImageFilter.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
itk_wrap_include(itkCudaImage.h)
itk_wrap_include(itkCudaImageToImageFilter.h)
itk_wrap_include(itkInPlaceImageFilter.h)

itk_wrap_class("itk::CudaImageToImageFilter" POINTER)

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}")

cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "" "")
Copy link
Collaborator

@SimonRit SimonRit Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this necessary? Which filter requires no parent? I think it's better to avoid this parent option and to make it mandatory

cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "IP" "itk::InPlaceImageFilter")

cuda_wrap_vector_combinations("${vector_types}" "" "")
Copy link
Collaborator

@SimonRit SimonRit Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Idem

cuda_wrap_vector_combinations("${vector_types}" "IP" "itk::InPlaceImageFilter")

itk_end_wrap_class()
15 changes: 15 additions & 0 deletions wrapping/itkCudaInPlaceImageFilter.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,15 @@
itk_wrap_include(itkCudaImage.h)

itk_wrap_class("itk::CudaInPlaceImageFilter" POINTER)

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}")

cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "IP")
Copy link
Collaborator

@SimonRit SimonRit Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Idem

cuda_wrap_image_filter_combinations("${from_types}" "${to_types}" "IP" "itk::InPlaceImageFilter")

cuda_wrap_vector_combinations("${vector_types}" "IP")
Copy link
Collaborator

@SimonRit SimonRit Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Idem

cuda_wrap_vector_combinations("${vector_types}" "IP" "itk::InPlaceImageFilter")

itk_end_wrap_class()
13 changes: 13 additions & 0 deletions wrapping/itkCudaSquareImageFilter.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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()
4 changes: 3 additions & 1 deletion wrapping/itkImageSourceCudaCommon.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}>")
Copy link
Collaborator

@SimonRit SimonRit Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Isn't it already wrapped by ITK?

endforeach()
endforeach()

Expand All @@ -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}>")
Copy link
Collaborator

@SimonRit SimonRit Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Idem

endforeach()
endforeach()
endforeach()
Expand Down
31 changes: 20 additions & 11 deletions wrapping/itkImageToImageFilterCudaCommon.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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()
Expand Down
33 changes: 33 additions & 0 deletions wrapping/itkInPlaceImageFilterCudaCommon.wrap
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,33 @@
itk_wrap_include(itkCudaImage.h)
itk_wrap_include(itkInPlaceImageFilter.h)

itk_wrap_class("itk::InPlaceImageFilter" POINTER)

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(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()
foreach(c ${ITK_WRAP_VECTOR_COMPONENTS})
foreach(vt ${vector_types})
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()

itk_end_wrap_class()