The expiration time for new job artifacts in CI/CD pipelines is now 30 days (GitLab default). Previously generated artifacts in already completed jobs will not be affected by the change. The latest artifacts for all jobs in the latest successful pipelines will be kept. More information: https://gitlab.lrz.de/help/user/admin_area/settings/continuous_integration.html#default-artifacts-expiration

Commit 36cb84fc authored by Declara Denis's avatar Declara Denis Committed by Christian Schulte zu Berge
Browse files

Fixed two bugs related to image flipping in cuda-confidence maps calculation

parent 0244d965
...@@ -53,6 +53,7 @@ namespace cuda { ...@@ -53,6 +53,7 @@ namespace cuda {
bool isUpsideDown; bool isUpsideDown;
int imageWidth; int imageWidth;
int imageHeight; int imageHeight;
int iterationCount;
float solutionResidualNorm; float solutionResidualNorm;
float systemCreationTime; float systemCreationTime;
float systemSolveTime; float systemSolveTime;
...@@ -68,6 +69,12 @@ namespace cuda { ...@@ -68,6 +69,12 @@ namespace cuda {
_gpuData->systemSolveTime = 0.0f; _gpuData->systemSolveTime = 0.0f;
} }
CudaConfidenceMapsSystemSolver::~CudaConfidenceMapsSystemSolver()
{
// This also frees all the memory reserved on the GPU
delete _gpuData;
}
void CudaConfidenceMapsSystemSolver::uploadImage(const unsigned char* imageData, int imageWidth, int imageHeight, void CudaConfidenceMapsSystemSolver::uploadImage(const unsigned char* imageData, int imageWidth, int imageHeight,
float gradientScaling, float alpha, float beta, float gamma, float gradientScaling, float alpha, float beta, float gamma,
bool useGPU, bool isUpsideDown) { bool useGPU, bool isUpsideDown) {
...@@ -95,7 +102,7 @@ namespace cuda { ...@@ -95,7 +102,7 @@ namespace cuda {
for (int y = 0; y < height; ++y) { for (int y = 0; y < height; ++y) {
for (int x = 0; x < width; ++x) { for (int x = 0; x < width; ++x) {
float value; float value;
if (!isUpsideDown) { if (isUpsideDown) {
value = y / (height - 1.0f); value = y / (height - 1.0f);
} }
else { else {
...@@ -108,6 +115,7 @@ namespace cuda { ...@@ -108,6 +115,7 @@ namespace cuda {
// Uplaod solution vector on the GPU // Uplaod solution vector on the GPU
_gpuData->x_d = _gpuData->x_h; _gpuData->x_d = _gpuData->x_h;
std::cout << "Reset!" << std::endl;
} }
void CudaConfidenceMapsSystemSolver::solve(int maximumIterations, float errorTolerance) { void CudaConfidenceMapsSystemSolver::solve(int maximumIterations, float errorTolerance) {
...@@ -119,6 +127,7 @@ namespace cuda { ...@@ -119,6 +127,7 @@ namespace cuda {
cusp::precond::diagonal<float, cusp::device_memory> M(_gpuData->L_d); cusp::precond::diagonal<float, cusp::device_memory> M(_gpuData->L_d);
cusp::krylov::cg(_gpuData->L_d, _gpuData->x_d, _gpuData->b_d, monitor, M); cusp::krylov::cg(_gpuData->L_d, _gpuData->x_d, _gpuData->b_d, monitor, M);
_gpuData->solutionResidualNorm = monitor.residual_norm(); _gpuData->solutionResidualNorm = monitor.residual_norm();
_gpuData->iterationCount = monitor.iteration_count();
// Downlaod the solution, which has been computed on the GPU to the CPU // Downlaod the solution, which has been computed on the GPU to the CPU
_gpuData->x_h = _gpuData->x_d; _gpuData->x_h = _gpuData->x_d;
...@@ -132,6 +141,11 @@ namespace cuda { ...@@ -132,6 +141,11 @@ namespace cuda {
return thrust::raw_pointer_cast(&_gpuData->x_h[0]); return thrust::raw_pointer_cast(&_gpuData->x_h[0]);
} }
int CudaConfidenceMapsSystemSolver::getSolutionIterationCount() const {
return _gpuData->iterationCount;
}
float CudaConfidenceMapsSystemSolver::getSolutionResidualNorm() const { float CudaConfidenceMapsSystemSolver::getSolutionResidualNorm() const {
return _gpuData->solutionResidualNorm; return _gpuData->solutionResidualNorm;
} }
...@@ -152,6 +166,7 @@ namespace cuda { ...@@ -152,6 +166,7 @@ namespace cuda {
// Resize the system vectors and matrices to accomodate the different image size // Resize the system vectors and matrices to accomodate the different image size
_gpuData->imageWidth = imageWidth; _gpuData->imageWidth = imageWidth;
_gpuData->imageHeight = imageHeight; _gpuData->imageHeight = imageHeight;
_gpuData->isUpsideDown = isUpsideDown;
int numElements = imageWidth * imageHeight; int numElements = imageWidth * imageHeight;
_gpuData->x_h.resize(numElements); _gpuData->x_h.resize(numElements);
_gpuData->b_d.resize(numElements); _gpuData->b_d.resize(numElements);
...@@ -301,7 +316,7 @@ namespace cuda { ...@@ -301,7 +316,7 @@ namespace cuda {
} }
static __global__ void k_buildSystem(float* L, int pitch, const unsigned char* image, int width, int height, static __global__ void k_buildSystem(float* L, int pitch, const unsigned char* image, int width, int height,
float gradientScaling, float alpha, float beta, float gamma) float gradientScaling, float alpha, float beta, float gamma, bool isUpsideDown)
{ {
const int x = blockDim.x * blockIdx.x + threadIdx.x; const int x = blockDim.x * blockIdx.x + threadIdx.x;
...@@ -313,11 +328,17 @@ namespace cuda { ...@@ -313,11 +328,17 @@ namespace cuda {
const int offsets[] = {-width-1, -width, -width+1, -1, 0, 1, width-1, width, width+1}; const int offsets[] = {-width-1, -width, -width+1, -1, 0, 1, width-1, width, width+1};
const float gammas[] = {gamma_sq2, 0.0f, gamma_sq2, gamma, 0.0f, gamma, gamma_sq2, 0.0f, gamma_sq2}; const float gammas[] = {gamma_sq2, 0.0f, gamma_sq2, gamma, 0.0f, gamma, gamma_sq2, 0.0f, gamma_sq2};
const float attenuations[] = { // Precompute the three attenuation values for the curernt pixel (the row above, current row, and row below)
1.0f - exp(-alpha * (1.0f - (y - 1.0f) / (height - 1.0f))), float attenuations[3];
1.0f - exp(-alpha * (1.0f - (y ) / (height - 1.0f))), if (isUpsideDown) {
1.0f - exp(-alpha * (1.0f - (y + 1.0f) / (height - 1.0f))) attenuations[0] = 1.0f - exp(-alpha * (1.0f - (y - 1.0f) / (height - 1.0f)));
}; attenuations[1] = 1.0f - exp(-alpha * (1.0f - (y ) / (height - 1.0f)));
attenuations[2] = 1.0f - exp(-alpha * (1.0f - (y + 1.0f) / (height - 1.0f)));
} else {
attenuations[0] = 1.0f - exp(-alpha * (y - 1.0f) / (height - 1.0f));
attenuations[1] = 1.0f - exp(-alpha * (y ) / (height - 1.0f));
attenuations[2] = 1.0f - exp(-alpha * (y + 1.0f) / (height - 1.0f));
}
// Filter off out-of-bounds edges // Filter off out-of-bounds edges
unsigned short filter = 495; // 111 101 111 unsigned short filter = 495; // 111 101 111
...@@ -373,7 +394,7 @@ namespace cuda { ...@@ -373,7 +394,7 @@ namespace cuda {
k_buildSystem<<<dimGrid, dimBlock>>>(thrust::raw_pointer_cast(&_gpuData->L_d.values.values[0]), _gpuData->L_d.values.pitch, k_buildSystem<<<dimGrid, dimBlock>>>(thrust::raw_pointer_cast(&_gpuData->L_d.values.values[0]), _gpuData->L_d.values.pitch,
thrust::raw_pointer_cast(&_gpuData->imageBuffer_d[0]), thrust::raw_pointer_cast(&_gpuData->imageBuffer_d[0]),
imageWidth, imageHeight, imageWidth, imageHeight,
gradientScaling, alpha, beta, gamma); gradientScaling, alpha, beta, gamma, _gpuData->isUpsideDown);
} }
} // cuda } // cuda
......
...@@ -10,6 +10,7 @@ namespace cuda { ...@@ -10,6 +10,7 @@ namespace cuda {
public: public:
CudaConfidenceMapsSystemSolver(); CudaConfidenceMapsSystemSolver();
~CudaConfidenceMapsSystemSolver();
/** /**
* Uploads an image to the solver. Additionally the matrices and vectors that are needed to * Uploads an image to the solver. Additionally the matrices and vectors that are needed to
...@@ -52,6 +53,12 @@ namespace cuda { ...@@ -52,6 +53,12 @@ namespace cuda {
*/ */
const float* getSolution(int& width, int& height); const float* getSolution(int& width, int& height);
/**
* Returns the number of CG iterations that were actually performed when solving the system
* \see solve() was called
*/
int getSolutionIterationCount() const;
/** /**
* Returns the residual norm of the solution as a measure of error; * Returns the residual norm of the solution as a measure of error;
*/ */
......
...@@ -40,7 +40,7 @@ namespace campvis { ...@@ -40,7 +40,7 @@ namespace campvis {
, p_inputImage("InputImage", "Input Image", "", DataNameProperty::READ) , p_inputImage("InputImage", "Input Image", "", DataNameProperty::READ)
, p_outputConfidenceMap("OutputConfidenceMap", "Output Confidence Map", "us.confidence", DataNameProperty::WRITE) , p_outputConfidenceMap("OutputConfidenceMap", "Output Confidence Map", "us.confidence", DataNameProperty::WRITE)
, p_resetResult("ResetSolution", "Reset solution vector") , p_resetResult("ResetSolution", "Reset solution vector")
, p_iterations("IterationCount", "Conjugate Gradient Iterations", 200, 1, 10000) , p_iterations("IterationCount", "Conjugate Gradient Iterations", 200, 1, 500)
, p_gradientScaling("GradientScaling", "Scaling factor for gradients", 2.0f, 0.001, 10) , p_gradientScaling("GradientScaling", "Scaling factor for gradients", 2.0f, 0.001, 10)
, p_paramAlpha("Alpha", "Alpha (TGC)", 2.0f, 0.001, 10) , p_paramAlpha("Alpha", "Alpha (TGC)", 2.0f, 0.001, 10)
, p_paramBeta("Beta", "Beta (Weight mapping)", 20.0f, 0.001, 200) , p_paramBeta("Beta", "Beta (Weight mapping)", 20.0f, 0.001, 200)
...@@ -107,7 +107,7 @@ namespace campvis { ...@@ -107,7 +107,7 @@ namespace campvis {
if (p_printStatistics.getValue()) { if (p_printStatistics.getValue()) {
std::cout << "Residual: " << _solver.getSolutionResidualNorm() << std::endl; std::cout << "Residual: " << _solver.getSolutionResidualNorm() << std::endl;
std::cout << "System Creation Time: " << _solver.getSystemCreationTime() << std::endl; std::cout << "System Creation Time: " << _solver.getSystemCreationTime() << std::endl;
std::cout << "System Solve Time: " << _solver.getSystemSolveTime() << std::endl; std::cout << "System Solve Time: " << _solver.getSystemSolveTime() << " (" << _solver.getSolutionIterationCount() << " iterations)" << std::endl;
} }
} }
} }
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment