2.12.2021, 9:00 - 11:00: Due to updates GitLab may be unavailable for some minutes between 09:00 and 11:00.

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

Integrated cuda code for confidence maps estimation.

Code still needs to be cleaned up a lot, and settings need to be exposed
to the GUI.
parent f2c44ee4
#include <chrono>
#include <iostream>
#include <iomanip>
#include "cm.h"
#include "cm_cusp.h"
void measureExecution(std::string message, const std::function<void()> &function)
{
std::cout << "started: " << message << std::endl;
auto start = std::chrono::high_resolution_clock::now();
function();
auto stop = std::chrono::high_resolution_clock::now();
float milliseconds = std::chrono::duration_cast<std::chrono::microseconds>(stop-start).count() / 1000.0f;
std::cout << "finished: " << std::left << std::setw(35) << message << "["
<< milliseconds << "ms]" << std::endl;
}
std::unique_ptr<ConfidenceMapSolver> createCUSPSolver(int width, int height, float gradientScaling,
float alpha, float beta, float gamma)
{
return std::make_unique<CuspConfidenceMapSolver>(width, height, gradientScaling, alpha, beta, gamma);
}
\ No newline at end of file
#pragma once
typedef unsigned char uint8;
#include <memory>
#include <vector>
#include <functional>
void measureExecution(std::string message, const std::function<void()> &function);
class ConfidenceMapSolver
{
public:
virtual ~ConfidenceMapSolver() {}
virtual void createSystem(const uint8* image, int width, int height) = 0;
virtual void setInitialSolution(const std::vector<float> &val) = 0;
virtual void solve() = 0;
virtual const float* getSolution(int &width, int &height) const = 0;
};
std::unique_ptr<ConfidenceMapSolver> createCUSPSolver(int width, int height, float gradientScaling,
float alpha, float beta, float gamma);
/*std::unique_ptr<ConfidenceMapSolver> createEIGENSolver(int width, int height, float gradientScaling,
float alpha, float beta, float gamma);*/
\ No newline at end of file
......@@ -11,7 +11,7 @@ CuspConfidenceMapSolver::~CuspConfidenceMapSolver()
CUSP_CM_destroyGpuData(gpuData);
}
void CuspConfidenceMapSolver::createSystem(const uint8* image, int width, int height)
void CuspConfidenceMapSolver::createSystem(const unsigned char* image, int width, int height)
{
if (width != this->width || height != this->height) {
CUSP_CM_resizeGpuData(*gpuData, width, height);
......@@ -30,15 +30,9 @@ void CuspConfidenceMapSolver::setInitialSolution(const std::vector<float> &val)
void CuspConfidenceMapSolver::solve()
{
measureExecution(" . uploading system", [=]() {
CUSP_CM_uploadSystem(*gpuData);
});
measureExecution(" . solving system", [=]() {
CUSP_CM_solveSystem(*gpuData, 100, 1e-20);
});
measureExecution(" . downloading solution", [=]() {
CUSP_CM_downloadSolution(*gpuData);
});
CUSP_CM_uploadSystem(*gpuData);
CUSP_CM_solveSystem(*gpuData, 100, 1e-20);
CUSP_CM_downloadSolution(*gpuData);
}
const float* CuspConfidenceMapSolver::getSolution(int &width, int &height) const
......
......@@ -60,25 +60,6 @@ void CUSP_CM_setInitialSolution(CuspGPUData& data, const std::vector<float>& val
void CUSP_CM_uploadSystem(CuspGPUData &data)
{
// Check for NaN values
/*
for (int r = 0; r < data.x_h.size(); ++r) {
for (int d = 0; d < 9; ++d) {
if (std::isnan(data.L_h.values(r, d))) {
std::cout << "NaN detected in A: " << r << std::endl;
break;
}
}
if (std::isnan(data.x_h[r])) {
std::cout << "NaN detected in x: " << r << std::endl;
}
if (std::isnan(data.b_h[r])) {
std::cout << "NaN detected in b: " << r << std::endl;
}
}*/
data.x_d = data.x_h;
data.b_d = data.b_h;
data.L_d = data.L_h;
......
#pragma once
#include "cm.h"
#include <vector>
class CuspGPUData;
/* Laplacian construction parameters */
class CuspConfidenceMapSolver : public ConfidenceMapSolver
class CuspConfidenceMapSolver
{
public:
CuspConfidenceMapSolver(int width, int height, float gradientScaling,
......@@ -13,7 +11,7 @@ public:
virtual ~CuspConfidenceMapSolver();
virtual void createSystem(const uint8* image, int width, int height);
virtual void createSystem(const unsigned char* image, int width, int height);
virtual void setInitialSolution(const std::vector<float> &val);
virtual void solve();
......@@ -24,8 +22,8 @@ private:
int width, height;
// Solver parameters
float alpha, beta, gamma;
float gradientScaling;
float alpha, beta, gamma;
// Matrices and Vectors
CuspGPUData* gpuData;
......
/*<
__global__ void k_buildCMSystem(cudaTextureObject_t image, int width, int height,
float alpha, float beta, float gamma,
float *out_weights)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;
if (x >= width || y >= height) return;
float attenuation_center = exp(-alpha * (float)y / height);
float attenuation_bottom = exp(-alpha * (float)(y+1) / height);
float c_center = tex2D(image, x, y) * attenuation_center;
float c_right = tex2D(image, x+1, y) * attenuation_center;
float c_bottom = tex2D(image, x, y+1) * attenuation_bottom;
float c_se = tex2D(image, x+1, y+1) * attenuation_bottom;
float c_sw = tex2D(image, x-1, y+1) * attenuation_bottom;
float w_right = exp(-beta * (abs(c_center - c_right) + gamma));
float w_bottom = exp(-beta * (abs(c_center - c_bottom)));
float w_se = exp(-beta * (abs(c_center - c_se) + sqrt(2.0f) * gamma));
float w_sw = exp(-beta * (abs(c_center - c_sw) + sqrt(2.0f) * gamma));
out_weights[idx*4 + 0] = w_right;
out_weights[idx*4 + 1] = w_sw;
out_weights[idx*4 + 2] = w_bottom;
out_weights[idx*4 + 3] = w_se;
}
int main()
{
// Create texture object
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypePitch2D
return 0;
}*/
\ No newline at end of file
......@@ -7,6 +7,7 @@ IF(${ModuleEnabled})
if(CUDA_FOUND)
# Source files:
FILE(GLOB ThisModSources RELATIVE ${ModulesDir}
modules/cudaconfidencemaps/core/*.cpp
modules/cudaconfidencemaps/pipelines/*.cpp
modules/cudaconfidencemaps/processors/*.cpp
)
......@@ -20,6 +21,15 @@ IF(${ModuleEnabled})
modules/cudaconfidencemaps/processors/*.h
)
# Build CUDA sources
set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-O3)
file(GLOB cuda_SOURCES modules/cudaconfidencemaps/core/*.cu)
cuda_add_library(
CudaConfidenceMaps_CUDA
${cuda_SOURCES}
)
LIST(APPEND ThisModExternalLibs CudaConfidenceMaps_CUDA)
SET(ThisModShaderDirectories "modules/cudaconfidencemaps/glsl")
SET(ThisModDependencies base io)
else()
......
......@@ -30,12 +30,12 @@ namespace campvis {
CudaConfidenceMapsDemo::CudaConfidenceMapsDemo(DataContainer* dc)
: AutoEvaluationPipeline(dc)
, _usReader()
, _usIgtlReader()
, _usBlurFilter(&_canvasSize)
//, _usResampler(&_canvasSize)
, _usMapsSolver()
{
addProcessor(&_usReader);
addProcessor(&_usIgtlReader);
addProcessor(&_usBlurFilter);
//addProcessor(&_usResampler);
addProcessor(&_usMapsSolver);
......@@ -49,11 +49,12 @@ namespace campvis {
AutoEvaluationPipeline::init();
// Create connectors
_usReader.p_url.setValue("/home/denis/projects/confidence_maps/build/data/gallbladder/392.png");
_usReader.p_importType.selectById("localIntensity");
_usReader.p_targetImageID.setValue("us.image");
_usReader.p_targetImageID.addSharedProperty(&_usBlurFilter.p_inputImage);
_usIgtlReader.p_receiveImages.setValue(true);
_usIgtlReader.p_receiveTransforms.setValue(false);
_usIgtlReader.p_receivePositions.setValue(false);
_usIgtlReader.p_targetImagePrefix.setValue("us.igtl.");
_usBlurFilter.p_inputImage.setValue("us.igtl.ImageClient");
_usBlurFilter.p_outputImage.setValue("us.blurred");
_usBlurFilter.p_outputImage.addSharedProperty(&_usMapsSolver.p_inputImage);
......
......@@ -31,6 +31,8 @@
#include "modules/preprocessing/processors/glgaussianfilter.h"
#include "modules/preprocessing/processors/glimageresampler.h"
#include "modules/cudaconfidencemaps/processors/cudaconfidencemapssolver.h"
#include "modules/openigtlink/processors/openigtlinkclient.h"
#include "core/properties/buttonproperty.h"
......@@ -63,6 +65,7 @@ namespace campvis {
//void execute();
protected:
OpenIGTLinkClient _usIgtlReader;
DevilImageReader _usReader;
GlGaussianFilter _usBlurFilter;
//GlImageResampler _usResampler;
......
......@@ -38,6 +38,7 @@ namespace campvis {
: AbstractProcessor()
, p_inputImage("InputImage", "Input Image", "", DataNameProperty::READ)
, p_outputConfidenceMap("OutputConfidenceMap", "Output Confidence Map", "us.confidence", DataNameProperty::WRITE)
, _solver(512, 512, 2.0f, 2.0f, 20.0f, 0.03f)
{
addProperty(p_inputImage);
......@@ -91,13 +92,16 @@ namespace campvis {
ImageRepresentationLocal::ScopedRepresentation img(data, p_inputImage.getValue());
if (img != 0) {
cgt::ivec3 size = img->getSize();
_solver.createSystem((unsigned char*)img->getWeaklyTypedPointer()._pointer, size.x, size.y);
_solver.solve();
const float *solution = _solver.getSolution(size.x, size.y);
ImageData *id = new ImageData(img->getParent()->getDimensionality(), size, img->getParent()->getNumChannels());
size_t elementCount = cgt::hmul(size);
uint8_t *imgData = new uint8_t[elementCount];
std::cout << "imgcount: " << elementCount << std::endl;
for (size_t i = 0; i < elementCount; ++i)
imgData[i] = (uint8_t)i;
WeaklyTypedPointer wtpData(WeaklyTypedPointer::UINT8, 1, imgData);
float *imgData = new float[elementCount];
memcpy(imgData, solution, sizeof(float)*elementCount);
WeaklyTypedPointer wtpData(WeaklyTypedPointer::FLOAT, 1, imgData);
ImageRepresentationLocal::create(id, wtpData);
id->setMappingInformation(img->getParent()->getMappingInformation());
data.addData(p_outputConfidenceMap.getValue(), id);
......
......@@ -41,6 +41,8 @@
#include "core/datastructures/imagedata.h"
#include "modules/cudaconfidencemaps/core/cm_cusp.h"
namespace campvis {
/**
......@@ -93,6 +95,7 @@ namespace campvis {
protected:
CuspConfidenceMapSolver _solver;
static const std::string loggerCat_;
......
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