BUG: Avoid 32-bit index/offset overflow in CUDA filters for large images by axel-grc · Pull Request #989 · RTKConsortium/RTK

axel-grc · 2026-08-28T12:49:49Z

Use 64-bit types for pointer offsets and element counts in the CUDA filters so that reconstructions on large projection stacks or volumes do not overflowing signed 32-bit arithmetic (which caused illegal memory access / segmentation faults).

SimonRit

Nice thanks. Does it address the problem reported on the mailing list?

SimonRit · 2026-08-28T12:55:59Z

This should be size_t

SimonRit · 2026-08-28T12:57:25Z

+    auto * pinOffset = pin + static_cast<ptrdiff_t>(nPixelsPerProj) * projectionOffset;
+    auto * poutOffset = pout + static_cast<ptrdiff_t>(nPixelsPerProj) * projectionOffset;


Are these static_cast needed?

SimonRit · 2026-08-28T12:58:27Z

+                         pinOffset,
+                         poutOffset,


I don't think we need intermediate variables if there is no static_cast

SimonRit · 2026-08-28T13:40:09Z

      matrix:
        python3-minor-version: ${{ github.event_name == 'pull_request' && fromJSON('["11"]') || fromJSON('["9","10","11"]') }}
        manylinux-platform: ${{ github.event_name == 'pull_request' && fromJSON('["_2_28-x64"]') || fromJSON('["_2_28-x64","2014-x64"]') }}
-        cuda-version: ${{ github.event_name == 'pull_request' && fromJSON('["118","130"]') || fromJSON('["118","124","128","130"]') }}


Why this change in this commit?

SimonRit · 2026-08-28T13:43:07Z

+    // Use 64-bit offset to avoid overflowing the signed 32-bit pointer arithmetic
+    // with large projection stacks


I don't think we need this comment

SimonRit · 2026-08-28T13:48:32Z

  cublasCreate(&handle);

  const double alpha = -1.0;
+#if CUDA_VERSION < 12000


You want to test the CUBLAS version, no CUDA

axel-grc · 2026-08-28T13:52:30Z

Nice thanks. Does it address the problem reported on the mailing list?

Yes, I am doing some tests on jean-zay, because my gpu has not enough ram

axel-grc · 2026-09-04T12:15:38Z

The changes seems to solve the problem @SimonRit

SimonRit · 2026-09-04T15:02:26Z


  const float alpha = -1.0;
-#if CUDA_VERSION < 12000
+#if CUBLAS_VER_MAJOR < 12


I guess we have a good reason to question the compatibility with Cuda 11 now.

Let's keep a cleanup for later and keep it as is

SimonRit

A bit of polishing required but looks good to me

SimonRit · 2026-09-04T15:04:52Z


  // Compute the index of the initial voxel
-  long int id = (k * c_Size.y + j) * c_Size.x + i;
+  long int id = (static_cast<long int>(k) * c_Size.y + j) * c_Size.x + i;


Interesting, the long int id could have been an unsigned int before this I guess. BTW, wouldn't unsigned long int make more sense?

SimonRit · 2026-09-04T15:07:14Z

  unsigned int i = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
  unsigned int j = __umul24(blockIdx_y, blockDim.y) + threadIdx.y;
-  unsigned int k = __umul24(blockIdx_z, blockDim.z) + threadIdx.z;
+  size_t       k = __umul24(blockIdx_z, blockDim.z) + threadIdx.z;


We should decide which strategy is the best between this and static_cast and uniformize everything

SimonRit · 2026-09-04T15:08:34Z

Declare it as long int once and for all?

SimonRit · 2026-09-04T15:10:21Z

+      // Compute indices in the volume (using 64-bit arithmetic to avoid overflow on large volumes)
+      const size_t volSize_X = static_cast<size_t>(c_volSize.x);
+      const size_t volSize_XY = volSize_X * static_cast<size_t>(c_volSize.y);
+      const size_t pxl = static_cast<size_t>(pos_low.x);
+      const size_t pxh = static_cast<size_t>(pos_high.x);
+      const size_t pyl = static_cast<size_t>(pos_low.y);
+      const size_t pyh = static_cast<size_t>(pos_high.y);
+      const size_t pzl = static_cast<size_t>(pos_low.z);
+      const size_t pzh = static_cast<size_t>(pos_high.z);
+      indices[0] = static_cast<long int>(pxl + pyl * volSize_X + pzl * volSize_XY); // index000
+      indices[1] = static_cast<long int>(pxl + pyl * volSize_X + pzh * volSize_XY); // index001
+      indices[2] = static_cast<long int>(pxl + pyh * volSize_X + pzl * volSize_XY); // index010
+      indices[3] = static_cast<long int>(pxl + pyh * volSize_X + pzh * volSize_XY); // index011
+      indices[4] = static_cast<long int>(pxh + pyl * volSize_X + pzl * volSize_XY); // index100
+      indices[5] = static_cast<long int>(pxh + pyl * volSize_X + pzh * volSize_XY); // index101
+      indices[6] = static_cast<long int>(pxh + pyh * volSize_X + pzl * volSize_XY); // index110
+      indices[7] = static_cast<long int>(pxh + pyh * volSize_X + pzh * volSize_XY); // index111


Mix between size_t and long int is missing, the static_cast are not necessary for the indices I believe then

SimonRit

There are quite a lot of things that need to be addressed, see detailed comments. This is an opportunity to clean the code and the PR is not always going in the right direction... Consider conditionally defining SizeValueType3 and SizeValueType4 with int3/long3 and int4/long4.

SimonRit · 2026-09-11T13:57:53Z

  itk-module-deps: "CudaCommon@main"
+  warnings-to-ignore: |
+    "warning #1388-D"
+    "warning #1394-D"


That only solves the problem for the CI which is not adequate. I believe this has been handled previously by a trick in itkCudaUtilities.hcu. I suggest to include itkIntTypes.h in this block of rtkCudaUtilities.hcu and to include rtkCudaUtilities.hcu when you need to use ITK's int types.

SimonRit · 2026-09-11T13:58:50Z


  const float alpha = -1.0;
-#if CUDA_VERSION < 12000
+#if CUBLAS_VER_MAJOR < 12


Let's keep a cleanup for later and keep it as is

SimonRit · 2026-09-11T14:00:30Z


  // Reset output volume
-  size_t memorySizeOutput = size[0] * size[1] * size[2] * size[3] * sizeof(float);
+  size_t memorySizeOutput = static_cast<size_t>(size[0]) * size[1] * size[2] * size[3] * sizeof(float);


Idem, don't static_cast but redefine size

SimonRit · 2026-09-11T14:01:31Z

+// System-adaptive buffer index types (ITK)
+using SizeValueType = itk::SizeValueType;


SimonRit · 2026-09-11T14:02:18Z


  // Reset output volume
-  size_t memorySizeOutput = size[0] * size[1] * size[2] * sizeof(float);
+  size_t memorySizeOutput = static_cast<size_t>(size[0]) * size[1] * size[2] * sizeof(float);


Idem, redefine size

SimonRit · 2026-09-15T15:07:00Z

    ray.d = pixelPos - ray.o;

-    int projOffset = numThread + proj * c_projSize.x * c_projSize.y;
+    size_t projOffset = static_cast<size_t>(numThread) + static_cast<size_t>(proj) * c_projSize.x * c_projSize.y;


SizeValueType

SimonRit · 2026-09-15T15:07:11Z

+// System-adaptive buffer index types (ITK)
+using SizeValueType = itk::SizeValueType;
+using OffsetValueType = itk::OffsetValueType;
+


SimonRit · 2026-09-15T15:07:36Z

+  // Index row major into the volume (64-bit to avoid 32-bit overflow)
+  SizeValueType ii = i, jj = j, kk = k;
+  SizeValueType out_idx = ii + (jj + kk * out_dim.y) * out_dim.x;
+  SizeValueType current_idx;


SimonRit · 2026-09-15T15:08:10Z

+  size_t memorySizeOutput =
+    sizeof(float) * static_cast<size_t>(output_vol_dim[0]) * output_vol_dim[1] * output_vol_dim[2];


SizeValueType

SimonRit · 2026-09-15T15:08:29Z


  // CUDA device pointers
-  size_t nVoxelsOutput = inputSize.x * inputSize.y * inputSize.z;
+  size_t nVoxelsOutput = static_cast<size_t>(inputSize.x) * inputSize.y * inputSize.z;


SizeValueType and redefine inputSize

Buffer indices, offsets and element counts in the CUDA filters used 32-bit arithmetic, which overflowed on large volumes or projection stacks and caused illegal memory accesses. Widen them to the ITK system-adaptive types (itk::SizeValueType / itk::OffsetValueType, 64-bit with ITK_USE_64BITS_IDS), typed directly so that only genuinely required casts remain. Resolves the problem reported on the rtk-users mailing list: https://www.creatis.insa-lyon.fr/pipermail/rtk-users/2026-August/002215.html

axel-grc force-pushed the overflow branch from 78415c8 to ea045b6 Compare August 28, 2026 13:17

SimonRit requested changes Aug 28, 2026

View reviewed changes

axel-grc force-pushed the overflow branch 6 times, most recently from 1b7970c to b8a59c0 Compare September 4, 2026 12:12

axel-grc force-pushed the overflow branch from b8a59c0 to 84642f6 Compare September 4, 2026 13:06

SimonRit reviewed Sep 4, 2026

View reviewed changes

SimonRit requested changes Sep 4, 2026

View reviewed changes

axel-grc force-pushed the overflow branch 4 times, most recently from db4802d to 2f22bae Compare September 8, 2026 13:48

SimonRit added this to the RTK 3.0 milestone Sep 10, 2026

SimonRit self-requested a review September 10, 2026 09:01

SimonRit requested changes Sep 15, 2026

View reviewed changes

axel-grc force-pushed the overflow branch 2 times, most recently from a71aae0 to 147f918 Compare September 28, 2026 08:38

SimonRit self-requested a review September 28, 2026 09:32

axel-grc force-pushed the overflow branch from 147f918 to 67678ad Compare September 28, 2026 11:41

		auto * pinOffset = pin + static_cast<ptrdiff_t>(nPixelsPerProj) * projectionOffset;
		auto * poutOffset = pout + static_cast<ptrdiff_t>(nPixelsPerProj) * projectionOffset;

		// Use 64-bit offset to avoid overflowing the signed 32-bit pointer arithmetic
		// with large projection stacks

		// System-adaptive buffer index types (ITK)
		using SizeValueType = itk::SizeValueType;

		size_t memorySizeOutput =
		sizeof(float) * static_cast<size_t>(output_vol_dim[0]) * output_vol_dim[1] * output_vol_dim[2];

Conversation

axel-grc commented Aug 28, 2026

Uh oh!

SimonRit left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

axel-grc commented Aug 28, 2026

Uh oh!

axel-grc commented Sep 4, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

SimonRit left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

SimonRit left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants