Conversation
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
| #include "gpu/registration/globalregistration.h" | ||
| #include "gpu/registration/registrationtypes.h" | ||
| #include "gpu/registration/imageoperations.h" | ||
| #include "gpu/registration/imageoperations.h" |
There was a problem hiding this comment.
warning: duplicate include [readability-duplicate-include]
cpp/cmd/mrreggpu.cpp:28:
- #include "gpu/registration/imageoperations.h"
- #include "gpu/registration/imageoperations.h"
+ #include "gpu/registration/imageoperations.h"| File::Matrix::save_transform(registration_result.transformation, centre, matrix_filename); | ||
| } | ||
| if (!matrix_1tomid_filename.empty()) { | ||
| File::Matrix::save_transform(halfway_transforms->half, centre, matrix_1tomid_filename); |
There was a problem hiding this comment.
warning: unchecked access to optional value [bugprone-unchecked-optional-access]
File::Matrix::save_transform(halfway_transforms->half, centre, matrix_1tomid_filename);
^| File::Matrix::save_transform(halfway_transforms->half, centre, matrix_1tomid_filename); | ||
| } | ||
| if (!matrix_2tomid_filename.empty()) { | ||
| File::Matrix::save_transform(halfway_transforms->half_inverse, centre, matrix_2tomid_filename); |
There was a problem hiding this comment.
warning: unchecked access to optional value [bugprone-unchecked-optional-access]
File::Matrix::save_transform(halfway_transforms->half_inverse, centre, matrix_2tomid_filename);
^| if (!transformed_midway1_filenames.empty()) { | ||
| // Compute midpioint transforms in scanner space and then build a midway output header that can hold both images | ||
| using ProjectiveTransform = Eigen::Transform<default_type, 3, Eigen::Projective>; | ||
| const ProjectiveTransform half_projective(halfway_transforms->half_matrix); |
There was a problem hiding this comment.
warning: unchecked access to optional value [bugprone-unchecked-optional-access]
const ProjectiveTransform half_projective(halfway_transforms->half_matrix);
^| // Compute midpioint transforms in scanner space and then build a midway output header that can hold both images | ||
| using ProjectiveTransform = Eigen::Transform<default_type, 3, Eigen::Projective>; | ||
| const ProjectiveTransform half_projective(halfway_transforms->half_matrix); | ||
| const ProjectiveTransform half_inverse_projective(halfway_transforms->half_inverse_matrix); |
There was a problem hiding this comment.
warning: unchecked access to optional value [bugprone-unchecked-optional-access]
const ProjectiveTransform half_inverse_projective(halfway_transforms->half_inverse_matrix);
^| const Buffer<float> matrixBuffer = context.new_buffer_from_host_memory<float>(matrixData); | ||
|
|
||
| const MomentUniforms uniforms{ | ||
| .centre = {centreScanner.x(), centreScanner.y(), centreScanner.z(), 0.0f}, |
There was a problem hiding this comment.
warning: floating point literal has suffix 'f', which is not uppercase [readability-uppercase-literal-suffix]
| .centre = {centreScanner.x(), centreScanner.y(), centreScanner.z(), 0.0f}, | |
| .centre = {centreScanner.x(), centreScanner.y(), centreScanner.z(), 0.0F}, |
|
|
||
| std::array<float, kMomentCount> momentValues{}; | ||
| for (size_t i = 0; i < kMomentCount; ++i) { | ||
| std::memcpy(&momentValues[i], &momentBits[i], sizeof(float)); |
There was a problem hiding this comment.
warning: do not use array subscript when the index is not an integer constant expression [cppcoreguidelines-pro-bounds-constant-array-index]
std::memcpy(&momentValues[i], &momentBits[i], sizeof(float));
^|
|
||
| std::array<float, kMomentCount> momentValues{}; | ||
| for (size_t i = 0; i < kMomentCount; ++i) { | ||
| std::memcpy(&momentValues[i], &momentBits[i], sizeof(float)); |
There was a problem hiding this comment.
warning: do not use array subscript when the index is not an integer constant expression [cppcoreguidelines-pro-bounds-constant-array-index]
std::memcpy(&momentValues[i], &momentBits[i], sizeof(float));
^|
|
||
| context.dispatch_kernel(transformKernel, dispatch_grid); | ||
|
|
||
| return outputTexture; |
There was a problem hiding this comment.
warning: constness of 'outputTexture' prevents automatic move [performance-no-automatic-move]
return outputTexture;
^|
|
||
| context.dispatch_kernel(transformKernel, dispatch_grid); | ||
|
|
||
| return outputTexture; |
There was a problem hiding this comment.
warning: constness of 'outputTexture' prevents automatic move [performance-no-automatic-move]
return outputTexture;
^
daljit46
force-pushed
the
affine_reg_gpu
branch
from
e97165f to
cfb2dca
Compare
| #include <string> | ||
| #include <vector> | ||
| #include <utility> | ||
| #include <vector> |
There was a problem hiding this comment.
warning: duplicate include [readability-duplicate-include]
cpp/core/gpu/registration/initialisation.cpp:46:
- #include <utility>
- #include <vector>
+ #include <utility>| return false; | ||
| } | ||
|
|
||
| const Eigen::Vector3f values = solver.eigenvalues(); |
There was a problem hiding this comment.
warning: the const qualified variable 'values' is copy-constructed from a const reference; consider making it a const reference [performance-unnecessary-copy-initialization]
| const Eigen::Vector3f values = solver.eigenvalues(); | |
| const Eigen::Vector3f& values = solver.eigenvalues(); |
| } | ||
|
|
||
| const Eigen::Vector3f values = solver.eigenvalues(); | ||
| const Eigen::Matrix3f vectors = solver.eigenvectors(); |
There was a problem hiding this comment.
warning: the const qualified variable 'vectors' is copy-constructed from a const reference; consider making it a const reference [performance-unnecessary-copy-initialization]
| const Eigen::Matrix3f vectors = solver.eigenvectors(); | |
| const Eigen::Matrix3f& vectors = solver.eigenvectors(); |
| std::sort(indices.begin(), indices.end(), [&](int a, int b) { return values[a] > values[b]; }); | ||
|
|
||
| for (size_t i = 0; i < indices.size(); ++i) { | ||
| eigenvalues[static_cast<int>(i)] = values[indices[i]]; |
There was a problem hiding this comment.
warning: do not use array subscript when the index is not an integer constant expression [cppcoreguidelines-pro-bounds-constant-array-index]
eigenvalues[static_cast<int>(i)] = values[indices[i]];
^|
|
||
| for (size_t i = 0; i < indices.size(); ++i) { | ||
| eigenvalues[static_cast<int>(i)] = values[indices[i]]; | ||
| eigenvectors.col(static_cast<int>(i)) = vectors.col(indices[i]); |
There was a problem hiding this comment.
warning: do not use array subscript when the index is not an integer constant expression [cppcoreguidelines-pro-bounds-constant-array-index]
eigenvectors.col(static_cast<int>(i)) = vectors.col(indices[i]);
^| .transformation_matrix = {}, | ||
| }; | ||
| m_compute_context->write_object_to_buffer(m_joint_histogram_uniforms_buffer, joint_histogram_uniforms); | ||
| const uint32_t jointHistogramPartialsSize = (m_num_bins * m_num_bins) * m_joint_histogram_dispatch_grid.workgroup_count(); |
There was a problem hiding this comment.
warning: Value stored to 'jointHistogramPartialsSize' during its initialization is never read [clang-analyzer-deadcode.DeadStores]
const uint32_t jointHistogramPartialsSize = (m_num_bins * m_num_bins) * m_joint_histogram_dispatch_grid.workgroup_count();
^Additional context
cpp/core/gpu/registration/nmicalculator.cpp:198: Value stored to 'jointHistogramPartialsSize' during its initialization is never read
const uint32_t jointHistogramPartialsSize = (m_num_bins * m_num_bins) * m_joint_histogram_dispatch_grid.workgroup_count();
^| m_compute_context->clear_buffer(m_joint_histogram_mass_buffer); | ||
|
|
||
| assert(transformation.param_count() == m_degrees_of_freedom); | ||
| const auto moving_dispatch_grid = DispatchGrid::element_wise_texture(m_moving, m_min_max_moving_kernel.workgroup_size); |
There was a problem hiding this comment.
warning: Value stored to 'moving_dispatch_grid' during its initialization is never read [clang-analyzer-deadcode.DeadStores]
const auto moving_dispatch_grid = DispatchGrid::element_wise_texture(m_moving, m_min_max_moving_kernel.workgroup_size);
^Additional context
cpp/core/gpu/registration/nmicalculator.cpp:287: Value stored to 'moving_dispatch_grid' during its initialization is never read
const auto moving_dispatch_grid = DispatchGrid::element_wise_texture(m_moving, m_min_max_moving_kernel.workgroup_size);
^| m_compute_context->dispatch_kernel(m_joint_histogram_smooth_kernel, smooth_grid); | ||
|
|
||
| const uint32_t histogramSize = m_num_bins * m_num_bins; | ||
| const DispatchGrid merge_grid{.x = histogramSize}; |
There was a problem hiding this comment.
warning: Value stored to 'merge_grid' during its initialization is never read [clang-analyzer-deadcode.DeadStores]
const DispatchGrid merge_grid{.x = histogramSize};
^Additional context
cpp/core/gpu/registration/nmicalculator.cpp:311: Value stored to 'merge_grid' during its initialization is never read
const DispatchGrid merge_grid{.x = histogramSize};
^|
|
||
| if (m_output == CalculatorOutput::CostAndGradients) { | ||
| if (transformation.is_affine()) { | ||
| std::array<float, 12> params; |
There was a problem hiding this comment.
warning: uninitialized record type: 'params' [cppcoreguidelines-pro-type-member-init]
| std::array<float, 12> params; | |
| std::array<float, 12> params{}; |
|
|
||
| m_compute_context->write_object_to_buffer(m_gradients_uniforms_buffer, gradients_uniforms); | ||
| } else { | ||
| std::array<float, 6> params; |
There was a problem hiding this comment.
warning: uninitialized record type: 'params' [cppcoreguidelines-pro-type-member-init]
| std::array<float, 6> params; | |
| std::array<float, 6> params{}; |
1 participant
This work builds on top of #3238. It introduces a new C++ command called
mrreggpu(chosen randomly and without much thought) that performs affine registration of 3D images on the GPU using WebGPU compute shaders. The code is completely independent ofmrregister.It's not ready to be merged and not ready for review yet. It needs much refinement, but I'm posting this PR to gather early feedback.
The utility of this command is rather limited since it only performs affine registration on scalar images, with some other limitations. The primary aim, however, is to provide a first real-world example of the GPU compute API introduced in #3238 so that we have a reference on how to use the GPU in the codebase.
This is also intended as a stepping stone towards non-linear registration on the GPU. I spent some time experimenting with SVF/SyN-style deformation using the current GPU API and I think the approach is feasible, but I'd appreciate guidance on which direction would be most useful.
mrreggpucurrently supports:mrregister, with minor differences (not yet sure these are justified).mrregister, it doesn't register both images into an average space. Instead it registers in both directions and then uses Lie algebra averaging to compute the transform for the next step (see here.Notes on the current state
cpp/core/gpu/registrationandcpp/core/gpu/shaders.magic_enum(for enum <-> string conversion), but that should likely be split into a separate PR.I'm aware that this is a rather large PR, but any feedback would be welcome.