Grangeat-based 2D/3D image registration
Loading...
Searching...
No Matches
ProjectDRRCuboidMaskCPU.h
Go to the documentation of this file.
1#pragma once
2
3#include "Common.h"
4
5namespace reg23 {
6
20at::Tensor ProjectDRRCuboidMask_CPU(const at::Tensor &volumeSize, const at::Tensor &voxelSpacing,
21 const at::Tensor &homographyMatrixInverse, double sourceDistance,
22 int64_t outputWidth, int64_t outputHeight, const at::Tensor &outputOffset,
23 const at::Tensor &detectorSpacing);
24
29__host__ at::Tensor ProjectDRRCuboidMask_CUDA(const at::Tensor &volumeSize, const at::Tensor &voxelSpacing,
30 const at::Tensor &homographyMatrixInverse, double sourceDistance,
31 int64_t outputWidth, int64_t outputHeight, const at::Tensor &outputOffset,
32 const at::Tensor &detectorSpacing);
33
34
52template <typename T, std::size_t faceCount> __host__ __device__ T RayConvexPolyhedronDistance(
53 const std::array<Vec<T, 3>, faceCount> &facePoints, const std::array<Vec<T, 3>, faceCount> &faceOutUnitNormals,
55
56 constexpr T epsilon = 1e-12;
57
58 T entryLambda = -std::numeric_limits<T>::infinity();
59 T exitLambda = std::numeric_limits<T>::infinity();
60 // std::size_t entryIndex = 0;
61 for (std::size_t i = 0; i < faceCount; ++i) {
62 const T dot = VecDot(faceOutUnitNormals[i], rayUnitDirection);
63 const T numer = VecDot(faceOutUnitNormals[i], facePoints[i] - rayPoint);
64 // If dot is 0, the ray is parallel to the face; if dot is negative, increasing lambda moves you along the ray
65 // in the 'inward' direction relative to this face, and vice versa if dot is positive.
66 if (std::abs(dot) < epsilon) {
67 if (numer < -epsilon) return 0.0f;
68 continue;
69 }
70 T lambda = numer / dot;
71 if (dot > epsilon) {
72 // Moving in 'outward' direction, so could be exiting the polyhedron through this face
74 } else {
75 // Moving in 'inward' direction, so could be entering the polyhedron through this face
77 }
78 // early rejection
79 if (entryLambda > exitLambda) return 0.0f;
80 }
81 // `entryLambda` is now the largest value for which ray is moving 'into' face, and `entryIndex` is index of
82 // corresponding face. `exitLambda` is the smallest value for which ray is moving 'out of' face. This means that,
83 // **if the ray intersects with the polyhedron**, it will enter at `entryLambda` and leave at `exitLambda`.
84
85 return exitLambda > entryLambda ? exitLambda - entryLambda : 0.0f;
86
87 // const Vec<T, 3> centrePoint = rayPoint + 0.5 * (entryLambda + exitLambda) * rayUnitDirection;
88 // // Checking that the centre point between the entry and exit points lies 'within' all the other faces of the
89 // // polyhedron. If it lies 'outside' of any, it doesn't intersect the polyhedron. Technically, could just check the
90 // // entry point against only the faces adjacent to the intersected face, but that would require expensive analysis of
91 // // the face intersections.
92 // for (std::size_t i = 0; i < faceCount; ++i) {
93 // // if (i == entryIndex) continue;
94 // if (VecDot(centrePoint - facePoints[i], faceOutUnitNormals[i]) > epsilon) return 0.f;
95 // }
96 //
97 // return std::abs(exitLambda - entryLambda);
98}
99
100
108
113
124
125 __host__ static CommonData Common(const at::Tensor &volumeSize, const at::Tensor &voxelSpacing,
126 const at::Tensor &homographyMatrixInverse, double sourceDistance,
127 int64_t outputWidth, int64_t outputHeight, const at::Tensor &outputOffset,
128 const at::Tensor &detectorSpacing, at::DeviceType device) {
129 // volumeSize should be a 1D tensor of 3 longs on the chosen device
130 TORCH_CHECK(volumeSize.sizes() == at::IntArrayRef{3});
131 TORCH_CHECK(volumeSize.dtype() == at::kLong);
132 TORCH_INTERNAL_ASSERT(volumeSize.device().type() == device);
133 // voxelSpacing should be a 1D tensor of 3 doubles on the chosen device
134 TORCH_CHECK(voxelSpacing.sizes() == at::IntArrayRef{3});
135 TORCH_CHECK(voxelSpacing.dtype() == at::kDouble);
136 TORCH_INTERNAL_ASSERT(voxelSpacing.device().type() == device);
137 // homographyMatrixInverse should be of size (4, 4), contain doubles and be on the chosen device
138 TORCH_CHECK(homographyMatrixInverse.sizes() == at::IntArrayRef({4, 4}));
139 TORCH_CHECK(homographyMatrixInverse.dtype() == at::kDouble);
140 TORCH_INTERNAL_ASSERT(homographyMatrixInverse.device().type() == device);
141 // outputOffset should be a 1D tensor of 2 doubles on the chosen device
142 TORCH_CHECK(outputOffset.sizes() == at::IntArrayRef{2});
143 TORCH_CHECK(outputOffset.dtype() == at::kDouble);
144 TORCH_INTERNAL_ASSERT(outputOffset.device().type() == device);
145 // detectorSpacing should be a 1D tensor of 2 doubles on the chosen device
146 TORCH_CHECK(detectorSpacing.sizes() == at::IntArrayRef{2});
147 TORCH_CHECK(detectorSpacing.dtype() == at::kDouble);
148 TORCH_INTERNAL_ASSERT(detectorSpacing.device().type() == device);
149
150 CommonData ret{};
151 ret.homographyMatrixInverse = Vec<Vec<float, 4>, 4>::FromTensor2D(homographyMatrixInverse);
152 const Vec<float, 6> planeSigns = Vec<int, 6>{1, 1, 1, -1, -1, -1}.StaticCast<float>();
154 int64_t, 3>::FromTensor(volumeSize).StaticCast<float>();
155 ret.cuboidIn.facePoints = VecOuter(cuboidHalfSize, planeSigns);
156 ret.cuboidIn.faceOutUnitNormals = VecCat(Vec<Vec<float, 3>, 3>::Identity(),
157 -1.f * Vec<Vec<float, 3>, 3>::Identity());
158
160 ret.cuboidAbove.facePoints = VecOuter(aboveBelowHalfSize, planeSigns);
161 ret.cuboidBelow.facePoints = ret.cuboidAbove.facePoints;
162 const float zSum = aboveBelowHalfSize.Z() + cuboidHalfSize.Z();
163 for (Vec<float, 3> &v : ret.cuboidAbove.facePoints) v.Z() += zSum;
164 for (Vec<float, 3> &v : ret.cuboidBelow.facePoints) v.Z() -= zSum;
165 ret.cuboidAbove.faceOutUnitNormals = ret.cuboidIn.faceOutUnitNormals;
166 ret.cuboidBelow.faceOutUnitNormals = ret.cuboidIn.faceOutUnitNormals;
167
168 ret.sourcePositionTransformed = MatMul(ret.homographyMatrixInverse,
169 Vec<float, 4>{0.0f, 0.0f, static_cast<float>(sourceDistance), 1.0f}).
170 XYZ();
171 ret.outputOffset = Vec<float, 2>::FromTensor(outputOffset);
172 ret.detectorSpacing = Vec<float, 2>::FromTensor(detectorSpacing);
173 ret.flatOutput = torch::zeros(at::IntArrayRef({outputWidth * outputHeight}),
174 at::TensorOptions{at::Device{device}});
175 return ret;
176 }
177};
178
179} // namespace reg23
General tools and structs.
#define __host__
Definition Global.h:17
#define __device__
Definition Global.h:22
A simple vector class derived from std::array<T, N>, providing overrides for all useful operators.
Definition Vec.h:21
__host__ __device__ constexpr Vec< newT, N > StaticCast() const
Construct a Vec with the elements form this one cast to a new type.
Definition Vec.h:237
static __host__ Vec FromTensor(const at::Tensor &t)
Construct a Vec from a 1D PyTorch tensor.
Definition Vec.h:98
__host__ __device__ constexpr Vec< T,(Ns+...)> VecCat(const Vec< T, Ns > &... vecs)
reg23::Vec concatenation of any number of vectors
Definition Vec.h:841
__host__ __device__ constexpr T VecDot(const Vec< T, N > &lhs, const Vec< T, N > &rhs)
reg23::Vec dot product
Definition Vec.h:667
__host__ __device__ constexpr Vec< Vec< T, N1 >, N2 > VecOuter(const Vec< T, N1 > &lhs, const Vec< T, N2 > &rhs)
reg23::Vec outer product; returns a column major matrix of size N1 x N2.
Definition Vec.h:682
__host__ __device__ constexpr Vec< T, R > MatMul(const Vec< Vec< T, R >, C > &lhs, const Vec< T, C > &rhs)
Matrix-vector multiplication of the Vec struct.
Definition Vec.h:894
__host__ __device__ T RayConvexPolyhedronDistance(const std::array< Vec< T, 3 >, faceCount > &facePoints, const std::array< Vec< T, 3 >, faceCount > &faceOutUnitNormals, const Vec< T, 3 > &rayPoint, const Vec< T, 3 > &rayUnitDirection)
Calculate the length of the intersection between a ray and a convex polyhedron.
Definition ProjectDRRCuboidMaskCPU.h:52
at::Tensor ProjectDRRCuboidMask_CPU(const at::Tensor &volumeSize, const at::Tensor &voxelSpacing, const at::Tensor &homographyMatrixInverse, double sourceDistance, int64_t outputWidth, int64_t outputHeight, const at::Tensor &outputOffset, const at::Tensor &detectorSpacing)
Calculate the distances of the intersections between the DRR-generation rays and the domain of the vo...
Definition ProjectDRRCuboidMaskCPU.cpp:9
__host__ at::Tensor ProjectDRRCuboidMask_CUDA(const at::Tensor &volumeSize, const at::Tensor &voxelSpacing, const at::Tensor &homographyMatrixInverse, double sourceDistance, int64_t outputWidth, int64_t outputHeight, const at::Tensor &outputOffset, const at::Tensor &detectorSpacing)
An implementation of reg23::ProjectDRRCuboidMask_CPU that uses CUDA parallelisation.
Vec< TextureAddressMode, DIMENSIONALITY > StringsToAddressModes(const std::array< std::string_view, DIMENSIONALITY > &strings)
Definition Texture.h:44
Definition GridSample3DCPU.cpp:6
Definition ProjectDRRCuboidMaskCPU.h:114
at::Tensor flatOutput
Definition ProjectDRRCuboidMaskCPU.h:122
Vec< float, 2 > detectorSpacing
Definition ProjectDRRCuboidMaskCPU.h:117
Vec< Vec< float, 4 >, 4 > homographyMatrixInverse
Definition ProjectDRRCuboidMaskCPU.h:115
Cuboid cuboidAbove
Definition ProjectDRRCuboidMaskCPU.h:120
Cuboid cuboidIn
Definition ProjectDRRCuboidMaskCPU.h:119
Vec< float, 3 > sourcePositionTransformed
Definition ProjectDRRCuboidMaskCPU.h:118
Vec< float, 2 > outputOffset
Definition ProjectDRRCuboidMaskCPU.h:116
Cuboid cuboidBelow
Definition ProjectDRRCuboidMaskCPU.h:121
Definition ProjectDRRCuboidMaskCPU.h:109
Vec< Vec< float, 3 >, 6 > facePoints
Definition ProjectDRRCuboidMaskCPU.h:110
Vec< Vec< float, 3 >, 6 > faceOutUnitNormals
Definition ProjectDRRCuboidMaskCPU.h:111
Definition ProjectDRRCuboidMaskCPU.h:107
static __host__ CommonData Common(const at::Tensor &volumeSize, const at::Tensor &voxelSpacing, const at::Tensor &homographyMatrixInverse, double sourceDistance, int64_t outputWidth, int64_t outputHeight, const at::Tensor &outputOffset, const at::Tensor &detectorSpacing, at::DeviceType device)
Definition ProjectDRRCuboidMaskCPU.h:125