43 #define finite _finite
51 #define CHECK_VEC(v1, v2) \
52 for (size_t j = 0; j < v1.size(); ++j) { \
53 if (v1[j] != v2[j]) { \
55 std::cout << i << ' ' << j << ' ' << v1[j] << ' ' << v2[j] << '\n'; \
58 #define DEBUG_FUNCN(v, func, input, N) \
59 if (__debug_pba && v.IsValid()) { \
60 vector<float> buf(v.GetLength()), buf_(v.GetLength()); \
61 for (int i = 0; i < N; ++i) { \
64 ProgramCU::FinishWorkCUDA(); \
66 v.CopyToHost(&buf_[0]); \
67 CHECK_VEC(buf, buf_); \
69 v.CopyToHost(&buf[0]); \
72 std::cout << #func << " : " << i << " : " << different << '\n'; \
75 #define DEBUG_FUNC(v, func, input) DEBUG_FUNCN(v, func, input, 2)
95 if (sz < 1024) std::cout <<
"WARNING: CUDA is unlikely to be supported!\n";
96 return sz < 1024 ? 0 : sz;
100 if (
sizeof(
CameraT) != 16 *
sizeof(
float)) exit(0);
117 const int* point_idx,
const int* cam_idx) {
168 if (!success && previous_allocated) {
176 if (
__verbose_level) std::cout <<
"WARNING: try not storing original JC\n";
183 if (
__verbose_level) std::cout <<
"WARNING: try not storing transpose JC\n";
190 if (
__verbose_level) std::cout <<
"WARNING: switch to memory saving mode\n";
209 std::cout <<
"Warm up device with storage allocation...\n";
237 if (
__verbose_level) std::cout <<
"NOTE: not storing camera Jacobian\n";
251 #ifdef PBA_CUDA_ALLOCATE_MORE
264 size_t sz0 = 800 * 600 * 2 * 4 *
sizeof(float);
265 size_t szq = q > 0 ? (
sizeof(float) * (q + m) * 4) : 0;
266 size_t sz =
sizeof(float) * (258 + 9 * n + 33 * m + 7 * k) + sz0;
269 sz += p * 6 *
sizeof(float);
271 sz += (2 * (k + q) *
sizeof(
float));
272 if (sz > total)
return false;
276 if (sz + szn > total)
283 if (sz + szn > total)
290 if (sz + szn > total)
298 if (sz + szn > total) {
300 sz -= (16 * k *
sizeof(float));
308 if (ncam <= 1 || npt <= 1 || nproj <= 1) {
321 #ifdef PBA_CUDA_ALLOCATE_MORE
325 ncam =
std::max(ncam, ncam_reserved);
327 nproj =
std::max(nproj, nproj_reserved);
337 std::cout <<
"Reserving storage for ncam = " << ncam <<
"; npt = " << npt
338 <<
"; nproj = " << nproj <<
'\n';
351 while (x < sz) x <<= 1;
363 #define REPORT_ALLOCATION(NAME) \
364 if (__verbose_allocation && NAME.GetDataSize() > 1024) \
365 std::cout << (NAME.GetDataSize() > 1024 * 1024 \
366 ? NAME.GetDataSize() / 1024 / 1024 \
367 : NAME.GetDataSize() / 1024) \
368 << (NAME.GetDataSize() > 1024 * 1024 ? "MB" : "KB") \
369 << "\t allocated for " #NAME "\n";
371 #define ASSERT_ALLOCATION(NAME) \
373 std::cerr << "WARNING: failed to allocate " \
374 << (__verbose_allocation ? #NAME "; size = " : "") \
375 << (total_sz / 1024 / 1024) << "MB + " \
376 << (NAME.GetRequiredSize() / 1024 / 1024) << "MB\n"; \
379 total_sz += NAME.GetDataSize(); \
380 REPORT_ALLOCATION(NAME); \
383 #define CHECK_ALLOCATION(NAME) \
384 if (NAME.GetDataSize() == 0 && NAME.GetRequiredSize() > 0) { \
385 ClearPreviousError(); \
386 std::cerr << "WARNING: unable to allocate " #NAME ": " \
387 << (NAME.GetRequiredSize() / 1024 / 1024) << "MB\n"; \
389 total_sz += NAME.GetDataSize(); \
390 REPORT_ALLOCATION(NAME); \
393 #define ALLOCATE_REQUIRED_DATA(NAME, num, channels) \
395 success &= NAME.InitTexture(num, 1, channels); \
396 ASSERT_ALLOCATION(NAME); \
399 #define ALLOCATE_OPTIONAL_DATA(NAME, num, channels, option) \
401 option = NAME.InitTexture(num, 1, channels); \
402 CHECK_ALLOCATION(NAME); \
404 NAME.InitTexture(0, 0, 0); \
413 vector<int> qmap, qlist;
414 vector<float> qmapw, qlistw;
470 for (
int i = 1; i <=
_num_camera; ++i) cpi[i] = cpi[i - 1] + cpnum[i - 1];
471 vector<int> cptidx = cpi;
486 for (
int i = 0; i <
_num_imgpt; ++i) ridx[cpidx[i]] = i;
501 for (
int i = 0, last_point = -1; i <
_num_imgpt; ++i) {
503 while (last_point < pt) ppi[++last_point] = i;
510 int* imp = &projection_map[i * 2];
525 std::cout <<
"Memory for Motion/Structure/Jacobian:\t"
526 << (total_sz / 1024 / 1024) <<
"MB\n";
531 vector<int>& qlist) {
555 if (iq < 0)
continue;
556 if (iq == i)
continue;
557 int ip = temp[2 * iq];
567 }
else if (ip == -1) {
569 temp[2 * iq + 1] = i;
571 temp[2 * i + 1] = iq;
575 temp[2 * i + 1] = iq;
576 temp[2 * ip + 1] = i;
582 std::cout <<
"Error: incorrect constraints\n";
589 int inext = temp[2 * i + 1];
590 if (inext == -1)
continue;
594 qmap.push_back(inext);
601 vector<float>& qmapw,
602 vector<float>& qlistw) {
609 if (qi == -1)
continue;
614 qpnum[qi] += cpnum[i];
621 if (qi == -1)
continue;
632 float wi = sqrt(qpnum[qi] / qcnum[qi]) *
_weight_q;
635 qmapw[i * 2 + 1] = wr;
636 qlistw[cidx * 2] = wi;
637 qlistw[cidx * 2 + 1] = wr;
675 int incompatible_radial_distortion = 0;
682 std::nth_element(focals.begin(), focals.begin() +
_num_camera / 2,
684 float median_focal_length = focals[
_num_camera / 2];
686 float radial_factor = median_focal_length * median_focal_length * 4.0f;
698 incompatible_radial_distortion++;
715 incompatible_radial_distortion++;
723 if (incompatible_radial_distortion) {
724 std::cout <<
"WARNING: incompatible radial distortion input; reset to 0;\n";
730 const float dist_bound = 1.0f;
735 int bad_point_count = 0;
739 float* rz = cam->
m[2];
741 oz[i] = (rz[0] * x[0] + rz[1] * x[1] + rz[2] * x[2] + cam->
t[2]);
745 float ozr = oz[i] / cam->
t[2];
748 float px = cam->
f * (cam->
m[0][0] * x[0] + cam->
m[0][1] * x[1] +
749 cam->
m[0][2] * x[2] + cam->
t[0]);
750 float py = cam->
f * (cam->
m[1][0] * x[0] + cam->
m[1][1] * x[1] +
751 cam->
m[1][2] * x[2] + cam->
t[1]);
754 if ((checkx && px * oz[i] * mx < 0 &&
fabs(mx) > 64) ||
755 (!checkx &&
py * oz[i] * my < 0 &&
fabs(my) > 64)) {
757 std::cout <<
"Warning: proj of #" << cmidx
758 <<
" on the wrong side, oz = " << oz[i] <<
" ("
759 << (px / oz[i]) <<
',' << (
py / oz[i]) <<
") (" << mx
760 <<
',' << my <<
")\n";
768 cpdist1[cmidx] =
std::min(cpdist1[cmidx], oz[i]);
770 cpdist2[cmidx] =
std::max(cpdist2[cmidx], oz[i]);
780 std::cout <<
"Enable data normalization on degeneracy\n";
785 std::nth_element(oz.begin(), oz.begin() +
_num_imgpt / 2, oz.end());
787 float shift_min =
std::min(oz_median * 0.001f, 1.0f);
788 float dist_threshold = shift_min * 0.1f;
792 << oz_median <<
")\n";
797 }
else if ((cpdist1[i] < dist_threshold ||
798 cpdist2[i] > -dist_threshold)) {
799 float shift = shift_min;
804 cpdist1[i] < dist_threshold && cpdist2[i] > -dist_threshold;
807 std::cout <<
"Adjust C" << std::setw(5) << i <<
" by "
808 << std::setw(12) << shift <<
" [B" << std::setw(2)
809 << cambpj[i] <<
"/" << std::setw(5) << camnpj[i] <<
"] ["
810 << (boths ?
'X' :
' ') <<
"][" << cpdist1[i] <<
", "
811 << cpdist2[i] <<
"]\n";
827 <<
" points are behind cameras.\n";
830 <<
" camera moved to avoid degeneracy.\n";
901 double e1 = 0, e2 = 0;
911 float z = r[6] * p[0] + r[7] * p[1] + r[8] * p[2] + t[2];
912 float xx = (r[0] * p[0] + r[1] * p[1] + r[2] * p[2] + t[0]);
913 float yy = (r[3] * p[0] + r[4] * p[1] + r[5] * p[2] + t[1]);
917 float rn = (m[0] * m[0] + m[1] * m[1]) * c[13] + 1.0f;
918 dx1 = c[0] * x - m[0] * rn;
919 dy1 = c[0] * y - m[1] * rn;
920 e1 += (dx1 * dx1 + dy1 * dy1);
921 e2 += (dx1 * dx1 + dy1 * dy1) / (rn * rn);
923 float rn = (x * x + y * y) * c[13] + 1.0f;
924 dx1 = c[0] * x * rn - m[0];
925 dy1 = c[0] * y * rn - m[1];
926 e1 += (dx1 * dx1 + dy1 * dy1) / (rn * rn);
927 e2 += (dx1 * dx1 + dy1 * dy1);
929 dx1 = c[0] * x - m[0];
930 dy1 = c[0] * y - m[1];
931 e1 += (dx1 * dx1 + dy1 * dy1);
932 e2 += (dx1 * dx1 + dy1 * dy1);
934 if (!isfinite(dx1) || !isfinite(dy1))
935 std::cout <<
"x = " << xx <<
" y = " << yy <<
" z = " << z <<
'\n'
936 <<
"t0 = " << t[0] <<
" t1 = " << t[1] <<
" t2 = " << t[2]
937 <<
'\n' <<
"p0 = " << p[0] <<
" p1 = " << p[1]
938 <<
" p2 = " << p[2] <<
'\n';
942 std::cout <<
"DEBUG: mean squared error = " << e1
943 <<
" in undistorted domain;\n";
944 std::cout <<
"DEBUG: mean squared error = " << e2
945 <<
" in distorted domain.\n";
975 std::cout <<
"Memory for Conjugate Gradient Solver:\t"
976 << (total_sz / 1024 / 1024) <<
"MB\n";
1038 }
else if (mode == 2) {
1041 }
else if (mode == 1)
1063 float& g_norm,
float& g_inf) {
1147 if (use_diagonal_q) {
1223 float_t alpha0 = rtz0 / (qtq0 + lambda * pdp0 - uv0);
1226 std::cout <<
" --0,\t alpha = " << alpha0
1228 if (!isfinite(alpha0)) {
1245 float_t rtzk = rtz0, rtz_min = rtz0, betak;
1266 betak = rtzk / rtzp;
1279 float_t alphak = rtzk / (qtqk + lambda * pdpk - uvk);
1283 std::cout <<
" --" << iteration <<
",\t alpha= " << alphak
1284 <<
", rtzk/rtz0 = " << rtz_ratio
1345 float_t alpha0 = rtz0 / (qtq0 + lambda * ptdp0);
1348 std::cout <<
" --0,\t alpha = " << alpha0
1350 if (!isfinite(alpha0)) {
1369 float_t rtzk = rtz0, rtz_min = rtz0, betak;
1391 betak = rtzk / rtzp;
1403 float_t alphak = rtzk / (qtqk + lambda * ptdpk);
1407 std::cout <<
" --" << iteration <<
",\t alpha= " << alphak
1408 <<
", rtzk/rtz0 = " << rtz_ratio
1501 float dx_sqnorm,
float damping) {
1502 float expected_reduction;
1513 expected_reduction = damping * dq + dxtg;
1515 expected_reduction = damping * dx_sqnorm + dxtg;
1528 expected_reduction = damping * dq + dxtg;
1530 expected_reduction = damping * dx_sqnorm + dxtg;
1540 expected_reduction = 2.0f * dxtg - njx;
1542 if (expected_reduction <= 0)
1543 expected_reduction = 0.001f * residual_reduction;
1547 expected_reduction = damping * dq + dxtg;
1549 expected_reduction = damping * dx_sqnorm + dxtg;
1561 return float(residual_reduction / expected_reduction);
1591 float mse_convert_ratio =
1607 std::cout <<
"Initial " << (
__verbose_sse ?
"sumed" :
"mean")
1608 <<
" squared error = " <<
__initial_mse * error_display_ratio
1609 <<
"\n----------------------------------------------\n";
1622 std::cout << std::left;
1629 if (num_cg_iteration == 0) {
1631 std::cout <<
"#" << std::setw(3) << i <<
" quit on numeric errors\n";
1638 std::cout <<
"#" << std::setw(3) << i <<
" 0 I e=" << std::setw(edwidth)
1640 <<
" u=" << std::setprecision(3) << std::setw(9) << damping
1641 <<
'\n' << std::setprecision(6);
1643 damping = damping * damping_adjust;
1644 damping_adjust = 2.0f * damping_adjust;
1660 std::cout <<
"#" << std::setw(3) << i <<
" " << std::setw(3)
1662 <<
" quit on too small change (" << dx_norm <<
" < "
1670 float average_residual = new_residual * mse_convert_ratio;
1674 if (isfinite(new_residual) && residual_reduction > 0) {
1693 std::cout <<
"#" << std::setw(3) << i <<
" " << std::setw(3)
1695 <<
" e=" << std::setw(edwidth)
1696 << average_residual * error_display_ratio
1697 <<
" u=" << std::setprecision(3) << std::setw(9) << damping
1698 <<
" r=" << std::setw(6)
1699 <<
floor(gain_ratio * 1000.f) * 0.001f
1700 <<
" g=" << std::setw(g_norm > 0 ? 9 : 1) << g_norm <<
" "
1701 << std::setw(9) << relative_reduction <<
' ' << std::setw(9)
1703 << std::setprecision(6);
1708 std::cout <<
"#" << std::setw(3) << i <<
" used up time budget.\n";
1713 std::cout <<
"#" << std::setw(3) << i
1714 <<
" converged with small gradient\n";
1719 std::cout <<
"#" << std::setw(3) << i <<
" satisfies MSE threshold\n";
1724 float temp = gain_ratio * 2.0f - 1.0f;
1725 float adaptive_adjust = 1.0f - temp * temp * temp;
1726 float auto_adjust =
std::max(1.0f / 3.0f, adaptive_adjust);
1729 damping = damping * auto_adjust;
1730 damping_adjust = 2.0f;
1742 std::cout <<
"#" << std::setw(3) << i <<
" " << std::setw(3)
1744 <<
" e=" << std::setw(edwidth) << std::left
1745 << average_residual * error_display_ratio
1746 <<
" u=" << std::setprecision(3) << std::setw(9) << damping
1750 <<
" --------- " << std::setw(9) << dx_norm
1752 << std::setprecision(6);
1758 damping_adjust = 2.0f;
1760 std::cout <<
"NOTE: switch to damping with an identity matix\n";
1763 damping = damping * damping_adjust;
1764 damping_adjust = 2.0f * damping_adjust;
1779 #define PROFILE_(A, B) \
1780 BundleTimerStart(TIMER_PROFILE_STEP); \
1781 for (int i = 0; i < repeat; ++i) { \
1785 BundleTimerSwitch(TIMER_PROFILE_STEP); \
1786 std::cout << std::setw(24) << A << ": " \
1787 << (BundleTimerGet(TIMER_PROFILE_STEP) / repeat) << "\n";
1789 #define PROFILE(A, B) PROFILE_(#A, A B)
1790 #define PROXILE(A, B) PROFILE_(A, B)
1794 std::cout <<
"---------------------------------\n"
1795 "| Run profiling steps ("
1796 << repeat <<
") |\n"
1797 "---------------------------------\n"
1802 PROXILE(
"Upload Measurements",
1806 std::cout <<
"---------------------------------\n";
1823 <<
"---------------------------------\n";
1838 std::cout <<
"---------------------------------\n";
1849 std::cout <<
"---------------------------------\n";
1861 std::cout <<
"---------------------------------\n";
1864 std::cout <<
"---------------------------------\n";
1884 std::cout <<
"---------------------------------\n"
1885 "| Not storing original JC | \n"
1886 "---------------------------------\n";
1893 std::cout <<
"---------------------------------\n"
1894 "| Not storing transpose JC | \n"
1895 "---------------------------------\n";
1907 std::cout <<
"---------------------------------\n"
1908 "| Not storing original JC | \n"
1909 "---------------------------------\n";
1915 std::cout <<
"---------------------------------\n"
1916 "| Not storing Camera Jacobians | \n"
1917 "---------------------------------\n";
1929 std::cout <<
"---------------------------------\n"
1930 "| Not storing any jacobians |\n"
1931 "---------------------------------\n";
1938 std::cout <<
"---------------------------------\n";
1951 ofstream out1(
"../../matlab/cg_j.txt");
1952 ofstream out2(
"../../matlab/cg_b.txt");
1953 ofstream out3(
"../../matlab/cg_x.txt");
1955 out1 << std::setprecision(20);
1956 out2 << std::setprecision(20);
1957 out3 << std::setprecision(20);
1963 vector<float> dx(plen);
1971 out2 << ee[i * 2] <<
' ' << ee[i * 2 + 1] <<
' ';
1973 float *cp = &jc[i * 16], *pp = &jp[i * 8];
1974 int cmin = cidx * 8, pmin = 8 *
_num_camera + pidx * 4;
1975 for (
int j = 0; j < 8; ++j)
1976 out1 << (i * 2 + 1) <<
' ' << (cmin + j + 1) <<
' ' << cp[j] <<
'\n';
1977 for (
int j = 0; j < 8; ++j)
1978 out1 << (i * 2 + 2) <<
' ' << (cmin + j + 1) <<
' ' << cp[j + 8] <<
'\n';
1979 for (
int j = 0; j < 4; ++j)
1980 out1 << (i * 2 + 1) <<
' ' << (pmin + j + 1) <<
' ' << pp[j] <<
'\n';
1981 for (
int j = 0; j < 4; ++j)
1982 out1 << (i * 2 + 2) <<
' ' << (pmin + j + 1) <<
' ' << pp[j + 4] <<
'\n';
1985 for (
size_t i = 0; i < dx.size(); ++i) out3 << dx[i] <<
' ';
1987 std::cout <<
"lambda = " << std::setprecision(20) << lambda <<
'\n';
#define ALLOCATE_REQUIRED_DATA(NAME, num, channels)
#define DEBUG_FUNCN(v, func, input, N)
#define ALLOCATE_OPTIONAL_DATA(NAME, num, channels, option)
bool __jacobian_normalize
bool __verbose_cg_iteration
bool __jc_store_transpose
void PrintBundleStatistics()
int __num_camera_modified
void BundleTimerSwap(int timer1, int timer2)
bool __lm_use_diagonal_damp
bool __depth_degeneracy_fix
float __lm_delta_threshold
float __depth_check_epsilon
float __cg_norm_threshold
void SaveBundleStatistics(int ncam, int npt, int nproj)
int __use_radial_distortion
bool __cg_schur_complement
bool __accurate_gain_ratio
void SaveBundleRecord(int iter, float res, float damping, float gn, float gi)
float __data_normalize_median
float __lm_damping_auto_switch
float __lm_gradient_threshold
bool __reset_initial_distortion
int __cg_recalculate_freq
int __num_projection_eval
bool __multiply_jx_usenoj
float BundleTimerGetNow(int timer=TIMER_OPTIMIZATION)
void ResetTemporarySetting()
int __bundle_current_mode
void ResetBundleStatistics()
bool __save_gradient_norm
bool IsTimeBudgetAvailable()
void CopyFromHost(const void *buf)
size_t GetReservedWidth()
void SwapData(CuTexImage &src)
void CopyToHost(void *buf)
bool InitTexture(unsigned int width, unsigned int height, unsigned int nchannel=1)
void SetTexture(void *data, unsigned int width, unsigned int nchannel=1)
@ STATUS_MEASURMENT_MISSING
@ STATUS_PROJECTION_MISSING
CuTexImage _cuCameraMeasurementListT
void AdjustBundleAdjsutmentMode()
void NonlinearOptimizeLM()
CuTexImage _cuCameraQListW
float EvaluateProjection(CuTexImage &cam, CuTexImage &point, CuTexImage &proj)
bool CheckRequiredMem(int fresh=1)
void ComputeJtE(CuTexImage &E, CuTexImage &JtE, int mode=0)
CuTexImage _cuCameraMeasurementMap
virtual void SetCameraData(size_t ncam, CameraT *cams)
void EvaluateJacobians(bool shuffle=true)
virtual float GetMeanSquaredError()
CuTexImage _cuVectorJtE
LM normal equation.
virtual void SetPointData(size_t npoint, Point3D *pts)
virtual void SetFocalMask(const int *fmask, float weight)
CuTexImage _cuCameraDataEX
void SaveBundleRecord(int iter, float res, float damping, float &g_norm, float &g_inf)
bool InitializeBundleGPU()
virtual int RunBundleAdjustment()
int SolveNormalEquation(float lambda)
bool InitializeStorageForCG()
float SaveUpdatedSystem(float residual_reduction, float dx_sqnorm, float damping)
void TransferDataToHost()
void ComputeBlockPC(float lambda, bool dampd=true)
float UpdateCameraPoint(CuTexImage &dx, CuTexImage &cuImageTempProj)
const float * _imgpt_data
CuTexImage _cuJacobianPoint
float EvaluateProjectionX(CuTexImage &cam, CuTexImage &point, CuTexImage &proj)
void PrepareJacobianNormalization()
void ApplyBlockPC(CuTexImage &v, CuTexImage &pv, int mode=0)
CuTexImage _cuJacobianCamera
void RunTestIterationLM(bool reduced)
SparseBundleCU(int device)
std::vector< float > _imgpt_datax
void ReleaseAllocatedData()
void ComputeJX(CuTexImage &X, CuTexImage &JX, int mode=0)
CuTexImage _cuCameraMeasurementList
void ProcessWeightCameraQ(std::vector< int > &cpnum, std::vector< int > &qmap, std::vector< float > &qmapw, std::vector< float > &qlistw)
void ComputeDiagonal(CuTexImage &JJ, CuTexImage &JJI)
virtual void SetProjection(size_t nproj, const Point2D *imgpts, const int *point_idx, const int *cam_idx)
bool ProcessIndexCameraQ(std::vector< int > &qmap, std::vector< int > &qlist)
CuTexImage _cuCameraQMapW
void SaveNormalEquation(float lambda)
CuTexImage _cuPointMeasurementMap
int SolveNormalEquationPCGX(float lambda)
CuTexImage _cuJacobianCameraT
CuTexImage _cuCameraQList
CuTexImage _cuMeasurements
CuTexImage _cuProjectionMap
int SolveNormalEquationPCGB(float lambda)
float EvaluateDeltaNorm()
void ReserveStorage(size_t ncam, size_t npt, size_t nproj)
CuTexImage _cuPointDataEX
void ReserveStorageAuto()
__host__ __device__ float2 fabs(float2 v)
static void error(char *msg)
MiniVec< float, N > floor(const MiniVec< float, N > &a)
void ComputeSXYPZ(float a, CuTexImage &texX, CuTexImage &texY, CuTexImage &texZ, CuTexImage &result)
void MultiplyBlockConditioner(int ncam, int npoint, CuTexImage &blocks, CuTexImage &vector, CuTexImage &result, int radial, int mode=0)
void ComputeDiagonalBlock_(float lambda, bool dampd, CuTexImage &camera, CuTexImage &point, CuTexImage &meas, CuTexImage &cmap, CuTexImage &cmlist, CuTexImage &pmap, CuTexImage &jmap, CuTexImage &jp, CuTexImage &sj, CuTexImage &diag, CuTexImage &blocks, bool intrinsic_fixed, int radial_distortion, bool add_existing_diagc, int mode=0)
void ComputeSQRT(CuTexImage &tex)
void ComputeJQtEC(CuTexImage &pe, CuTexImage &qlist, CuTexImage &wq, CuTexImage &sj, CuTexImage &result)
void ComputeJX_(CuTexImage &x, CuTexImage &jx, CuTexImage &camera, CuTexImage &point, CuTexImage &meas, CuTexImage &pjmap, bool intrinsic_fixed, int radial_distortion, int mode=0)
void ComputeSAX(float a, CuTexImage &texX, CuTexImage &result)
void ComputeVXY(CuTexImage &texX, CuTexImage &texY, CuTexImage &result, unsigned int part=0, unsigned int skip=0)
double ComputeVectorNormW(CuTexImage &vector, CuTexImage &weight, CuTexImage &buf)
void ComputeJX(int point_offset, CuTexImage &x, CuTexImage &jc, CuTexImage &jp, CuTexImage &jmap, CuTexImage &result, int mode=0)
void ComputeProjectionQ(CuTexImage &camera, CuTexImage &qmap, CuTexImage &qw, CuTexImage &proj, int offset)
void ClearPreviousError()
void ComputeDiagonalBlock(float lambda, bool dampd, CuTexImage &jc, CuTexImage &cmap, CuTexImage &jp, CuTexImage &pmap, CuTexImage &cmlist, CuTexImage &diag, CuTexImage &blocks, int radial_distortion, bool jc_transpose, bool add_existing_diagc, int mode=0)
void ComputeRSQRT(CuTexImage &tex)
size_t GetCudaMemoryCap()
void ComputeProjectionX(CuTexImage &camera, CuTexImage &point, CuTexImage &meas, CuTexImage &proj_map, CuTexImage &proj, int radial)
void ComputeSAXPY(float a, CuTexImage &texX, CuTexImage &texY, CuTexImage &result)
int SetCudaDevice(int device)
void ClearTextureObjectCache()
void ComputeDiagonalQ(CuTexImage &qlistw, CuTexImage &sj, CuTexImage &diag)
void ComputeJtE_(CuTexImage &e, CuTexImage &jte, CuTexImage &camera, CuTexImage &point, CuTexImage &meas, CuTexImage &cmap, CuTexImage &cmlist, CuTexImage &pmap, CuTexImage &jmap, CuTexImage &jp, bool intrinsic_fixed, int radial_distortion, int mode=0)
double ComputeVectorDot(CuTexImage &vector1, CuTexImage &vector2, CuTexImage &buf)
void ComputeProjection(CuTexImage &camera, CuTexImage &point, CuTexImage &meas, CuTexImage &proj_map, CuTexImage &proj, int radial)
float ComputeVectorMax(CuTexImage &vector, CuTexImage &buf)
void ComputeJQX(CuTexImage &x, CuTexImage &qmap, CuTexImage &wq, CuTexImage &sj, CuTexImage &jx, int offset)
void ComputeJacobian(CuTexImage &camera, CuTexImage &point, CuTexImage &jc, CuTexImage &jp, CuTexImage &proj_map, CuTexImage &sj, CuTexImage &meas, CuTexImage &cmlist, bool intrinsic_fixed, int radial_distortion, bool shuffle)
void ComputeDiagonal(CuTexImage &jc, CuTexImage &cmap, CuTexImage &jp, CuTexImage &pmap, CuTexImage &cmlist, CuTexImage &jtjd, CuTexImage &jtjdi, bool jc_transpose, int radial, bool add_existing_diagc)
void UpdateCameraPoint(int ncam, CuTexImage &camera, CuTexImage &point, CuTexImage &delta, CuTexImage &new_camera, CuTexImage &new_point, int mode=0)
void ComputeJtE(CuTexImage &pe, CuTexImage &jc, CuTexImage &cmap, CuTexImage &cmlist, CuTexImage &jp, CuTexImage &pmap, CuTexImage &jte, bool jc_transpose, int mode=0)
double ComputeVectorNorm(CuTexImage &vector, CuTexImage &buf)
bool ShuffleCameraJacobian(CuTexImage &jc, CuTexImage &map, CuTexImage &result)
void ResetCurrentDevice()
static size_t upgrade_dimension(size_t sz)