BUG: Avoid 32-bit index/offset overflow in CUDA filters for large images - #989
BUG: Avoid 32-bit index/offset overflow in CUDA filters for large images#989axel-grc wants to merge 1 commit into
Conversation
SimonRit
left a comment
SimonRit
left a comment
There was a problem hiding this comment.
Nice thanks. Does it address the problem reported on the mailing list?
| auto * pinOffset = pin + static_cast<ptrdiff_t>(nPixelsPerProj) * projectionOffset; | ||
| auto * poutOffset = pout + static_cast<ptrdiff_t>(nPixelsPerProj) * projectionOffset; |
There was a problem hiding this comment.
Are these static_cast needed?
| pinOffset, | ||
| poutOffset, |
There was a problem hiding this comment.
I don't think we need intermediate variables if there is no static_cast
| 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"]') }} |
There was a problem hiding this comment.
Why this change in this commit?
| // Use 64-bit offset to avoid overflowing the signed 32-bit pointer arithmetic | ||
| // with large projection stacks |
There was a problem hiding this comment.
I don't think we need this comment
| cublasCreate(&handle); | ||
|
|
||
| const double alpha = -1.0; | ||
| #if CUDA_VERSION < 12000 |
There was a problem hiding this comment.
You want to test the CUBLAS version, no CUDA
Yes, I am doing some tests on jean-zay, because my gpu has not enough ram |
1b7970c to
b8a59c0
Compare
|
The changes seems to solve the problem @SimonRit |
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).
|
|
||
| const float alpha = -1.0; | ||
| #if CUDA_VERSION < 12000 | ||
| #if CUBLAS_VER_MAJOR < 12 |
There was a problem hiding this comment.
I guess we have a good reason to question the compatibility with Cuda 11 now.
SimonRit
left a comment
SimonRit
left a comment
There was a problem hiding this comment.
A bit of polishing required but looks good to me
|
|
||
| // 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; |
There was a problem hiding this comment.
Interesting, the long int id could have been an unsigned int before this I guess. BTW, wouldn't unsigned long int make more sense?
| 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; |
There was a problem hiding this comment.
We should decide which strategy is the best between this and static_cast and uniformize everything
| { | ||
| int i = blockIdx.x * blockDim.x + threadIdx.x; | ||
| int j = blockIdx.y * blockDim.y + threadIdx.y; | ||
| int k = blockIdx.z * blockDim.z + threadIdx.z; |
There was a problem hiding this comment.
Declare it as long int once and for all?
| // 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 |
There was a problem hiding this comment.
Mix between size_t and long int is missing, the static_cast are not necessary for the indices I believe then
2 participants
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).