Skip to content

Commit

Permalink
remove gpu config stuff
Browse files Browse the repository at this point in the history
  • Loading branch information
@mp3guy
mp3guy committed Apr 30, 2019
1 parent 34385b5 commit 20402e5
Show file tree
Hide file tree
Showing 5 changed files with 32 additions and 220 deletions.
16 changes: 4 additions & 12 deletions Core/src/Cuda/cudafuncs.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -76,9 +76,7 @@ void icpStep(const mat33& Rcurr,
DeviceArray<JtJJtrSE3> & out,
float * matrixA_host,
float * vectorB_host,
float * residual_host,
int threads,
int blocks);
float * residual_host);

void rgbStep(const DeviceArray2D<DataTerm> & corresImg,
const float & sigma,
Expand All @@ -91,9 +89,7 @@ void rgbStep(const DeviceArray2D<DataTerm> & corresImg,
DeviceArray<JtJJtrSE3> & sum,
DeviceArray<JtJJtrSE3> & out,
float * matrixA_host,
float * vectorB_host,
int threads,
int blocks);
float * vectorB_host);

void so3Step(const DeviceArray2D<unsigned char> & lastImage,
const DeviceArray2D<unsigned char> & nextImage,
Expand All @@ -104,9 +100,7 @@ void so3Step(const DeviceArray2D<unsigned char> & lastImage,
DeviceArray<JtJJtrSO3> & out,
float * matrixA_host,
float * vectorB_host,
float * residual_host,
int threads,
int blocks);
float * residual_host);

void computeRgbResidual(const float & minScale,
const DeviceArray2D<short> & dIdx,
Expand All @@ -121,9 +115,7 @@ void computeRgbResidual(const float & minScale,
const float3 & kt,
const mat33 & krkinv,
int & sigmaSum,
int & count,
int threads,
int blocks);
int & count);

void createVMap(const CameraModel& intr,
const DeviceArray2D<unsigned short> & depth,
Expand Down
44 changes: 24 additions & 20 deletions Core/src/Cuda/reduce.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -396,9 +396,7 @@ void icpStep(const mat33& Rcurr,
DeviceArray<JtJJtrSE3> & out,
float * matrixA_host,
float * vectorB_host,
float * residual_host,
int threads,
int blocks)
float * residual_host)
{
int cols = vmap_curr.cols ();
int rows = vmap_curr.rows () / 3;
Expand Down Expand Up @@ -428,9 +426,12 @@ void icpStep(const mat33& Rcurr,
icp.N = cols * rows;
icp.out = sum;

icpKernel<<<blocks, threads>>>(icp);
constexpr int reduceThreads = 256;
constexpr int reduceBlocks = 80;

reduceSum<<<1, MAX_THREADS>>>(sum, out, blocks);
icpKernel<<<reduceBlocks, reduceThreads>>>(icp);

reduceSum<<<1, MAX_THREADS>>>(sum, out, reduceBlocks);

cudaSafeCall(cudaGetLastError());
cudaSafeCall(cudaDeviceSynchronize());
Expand Down Expand Up @@ -600,9 +601,7 @@ void rgbStep(const DeviceArray2D<DataTerm> & corresImg,
DeviceArray<JtJJtrSE3> & sum,
DeviceArray<JtJJtrSE3> & out,
float * matrixA_host,
float * vectorB_host,
int threads,
int blocks)
float * vectorB_host)
{
RGBReduction rgb;

Expand All @@ -619,9 +618,12 @@ void rgbStep(const DeviceArray2D<DataTerm> & corresImg,
rgb.N = rgb.cols * rgb.rows;
rgb.out = sum;

rgbKernel<<<blocks, threads>>>(rgb);
constexpr int reduceThreads = 256;
constexpr int reduceBlocks = 80;

rgbKernel<<<reduceBlocks, reduceThreads>>>(rgb);

reduceSum<<<1, MAX_THREADS>>>(sum, out, blocks);
reduceSum<<<1, MAX_THREADS>>>(sum, out, reduceBlocks);

cudaSafeCall(cudaGetLastError());
cudaSafeCall(cudaDeviceSynchronize());
Expand Down Expand Up @@ -845,9 +847,7 @@ void computeRgbResidual(const float & minScale,
const float3 & kt,
const mat33 & krkinv,
int & sigmaSum,
int & count,
int threads,
int blocks)
int & count)
{
int cols = nextImage.cols ();
int rows = nextImage.rows ();
Expand Down Expand Up @@ -880,15 +880,18 @@ void computeRgbResidual(const float & minScale,
rgb.N = cols * rows;
rgb.out = sumResidual;

residualKernel<<<blocks, threads>>>(rgb);
constexpr int reduceThreads = 256;
constexpr int reduceBlocks = 80;

residualKernel<<<reduceBlocks, reduceThreads>>>(rgb);

int2 out_host = {0, 0};
int2 * out;

cudaMalloc(&out, sizeof(int2));
cudaMemcpy(out, &out_host, sizeof(int2), cudaMemcpyHostToDevice);

reduceSum<<<1, MAX_THREADS>>>(sumResidual, out, blocks);
reduceSum<<<1, MAX_THREADS>>>(sumResidual, out, reduceBlocks);

cudaSafeCall(cudaGetLastError());
cudaSafeCall(cudaDeviceSynchronize());
Expand Down Expand Up @@ -1054,9 +1057,7 @@ void so3Step(const DeviceArray2D<unsigned char> & lastImage,
DeviceArray<JtJJtrSO3> & out,
float * matrixA_host,
float * vectorB_host,
float * residual_host,
int threads,
int blocks)
float * residual_host)
{
int cols = nextImage.cols();
int rows = nextImage.rows();
Expand All @@ -1078,9 +1079,12 @@ void so3Step(const DeviceArray2D<unsigned char> & lastImage,

so3.out = sum;

so3Kernel<<<blocks, threads>>>(so3);
constexpr int reduceThreads = 256;
constexpr int reduceBlocks = 80;

so3Kernel<<<reduceBlocks, reduceThreads>>>(so3);

reduceSum<<<1, MAX_THREADS>>>(sum, out, blocks);
reduceSum<<<1, MAX_THREADS>>>(sum, out, reduceBlocks);

cudaSafeCall(cudaGetLastError());
cudaSafeCall(cudaDeviceSynchronize());
Expand Down
175 changes: 0 additions & 175 deletions Core/src/Utils/GPUConfig.h
View file

This file was deleted.

16 changes: 4 additions & 12 deletions Core/src/Utils/RGBDOdometry.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -340,9 +340,7 @@ void RGBDOdometry::getIncrementalTransformation(Eigen::Vector3f & trans,
outDataSO3,
jtj.data(),
jtr.data(),
&residual[0],
GPUConfig::getInstance().so3StepThreads,
GPUConfig::getInstance().so3StepBlocks);
&residual[0]);
TOCK("so3Step");

lastSO3Error = sqrt(residual[0]) / residual[1];
Expand Down Expand Up @@ -452,9 +450,7 @@ void RGBDOdometry::getIncrementalTransformation(Eigen::Vector3f & trans,
kt,
krkInv,
sigma,
rgbSize,
GPUConfig::getInstance().rgbResThreads,
GPUConfig::getInstance().rgbResBlocks);
rgbSize);
TOCK("computeRgbResidual");
}

Expand Down Expand Up @@ -506,9 +502,7 @@ void RGBDOdometry::getIncrementalTransformation(Eigen::Vector3f & trans,
outDataSE3,
A_icp.data(),
b_icp.data(),
&residual[0],
GPUConfig::getInstance().icpStepThreads,
GPUConfig::getInstance().icpStepBlocks);
&residual[0]);
TOCK("icpStep");
}

Expand All @@ -532,9 +526,7 @@ void RGBDOdometry::getIncrementalTransformation(Eigen::Vector3f & trans,
sumDataSE3,
outDataSE3,
A_rgbd.data(),
b_rgbd.data(),
GPUConfig::getInstance().rgbStepThreads,
GPUConfig::getInstance().rgbStepBlocks);
b_rgbd.data());
TOCK("rgbStep");
}

Expand Down
1 change: 0 additions & 1 deletion Core/src/Utils/RGBDOdometry.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,6 @@
#include "../GPUTexture.h"
#include "../Cuda/cudafuncs.cuh"
#include "OdometryProvider.h"
#include "GPUConfig.h"

#include <vector>
#include <vector_types.h>
Expand Down

0 comments on commit 20402e5

Please sign in to comment.