Movatterモバイル変換

[0]ホーム
URL:

Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly

 Sign up
Appearance settings

Commit5e86bcd

Browse files
committed
only use the uchar4 cudacache normals
1 parentb4158a0 commit5e86bcd

File tree

12 files changed

+214
-96
lines changed

12 files changed

+214
-96
lines changed

‎FriedLiver/Source/CUDACache.cpp‎

Lines changed: 13 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -49,11 +49,15 @@ void CUDACache::storeFrame(const float* d_depth, unsigned int inputDepthWidth, u
4949
}
5050
CUDAImageUtil::convertDepthFloatToCameraSpaceFloat4(d_helperCamPos, d_inputDepth, *(float4x4*)&m_inputIntrinsicsInv, inputDepthWidth, inputDepthHeight);
5151
CUDAImageUtil::resampleFloat4(frame.d_cameraposDownsampled, m_width, m_height, d_helperCamPos, inputDepthWidth, inputDepthHeight);
52-
52+
#ifdef CUDACACHE_UCHAR_NORMALS
53+
CUDAImageUtil::computeNormals(d_helperNormals, d_helperCamPos, inputDepthWidth, inputDepthHeight);
54+
CUDAImageUtil::resampleFloat4(d_helperCamPos, m_width, m_height, d_helperNormals, inputDepthWidth, inputDepthHeight);//just use the memory from helpercampos
55+
CUDAImageUtil::convertNormalsFloat4ToUCHAR4(frame.d_normalsDownsampledUCHAR4, d_helperCamPos, m_width, m_height);
56+
#else
5357
CUDAImageUtil::computeNormals(d_helperNormals, d_helperCamPos, inputDepthWidth, inputDepthHeight);
5458
CUDAImageUtil::resampleFloat4(frame.d_normalsDownsampled, m_width, m_height, d_helperNormals, inputDepthWidth, inputDepthHeight);
55-
56-
CUDAImageUtil::convertNormalsFloat4ToUCHAR4(frame.d_normalsDownsampledUCHAR4, frame.d_normalsDownsampled, m_width, m_height);
59+
//CUDAImageUtil::convertNormalsFloat4ToUCHAR4(frame.d_normalsDownsampledUCHAR4, frame.d_normalsDownsampled, m_width, m_height);
60+
#endif
5761

5862
CUDAImageUtil::resampleFloat(frame.d_depthDownsampled, m_width, m_height, d_inputDepth, inputDepthWidth, inputDepthHeight);
5963

@@ -89,6 +93,11 @@ void CUDACache::fuseDepthFrames(CUDACache* globalCache, const int* d_validImages
8993
globalFrame.d_depthDownsampled, tmpFrame.d_intensityDerivsDownsampled, globalFrame.d_normalsDownsampledUCHAR4);
9094

9195
CUDAImageUtil::convertDepthFloatToCameraSpaceFloat4(globalFrame.d_cameraposDownsampled, globalFrame.d_depthDownsampled,MatrixConversion::toCUDA(m_intrinsicsInv), m_width, m_height);
96+
#ifdef CUDACACHE_UCHAR_NORMALS
97+
CUDAImageUtil::computeNormals(d_helperNormals, globalFrame.d_cameraposDownsampled, m_width, m_height);
98+
CUDAImageUtil::convertNormalsFloat4ToUCHAR4(globalFrame.d_normalsDownsampledUCHAR4, d_helperNormals, m_width, m_height);
99+
#else
92100
CUDAImageUtil::computeNormals(globalFrame.d_normalsDownsampled, globalFrame.d_cameraposDownsampled, m_width, m_height);
93-
CUDAImageUtil::convertNormalsFloat4ToUCHAR4(globalFrame.d_normalsDownsampledUCHAR4, globalFrame.d_normalsDownsampled, m_width, m_height);
101+
//CUDAImageUtil::convertNormalsFloat4ToUCHAR4(globalFrame.d_normalsDownsampledUCHAR4, globalFrame.d_normalsDownsampled, m_width, m_height);
102+
#endif
94103
}

‎FriedLiver/Source/CUDACache.h‎

Lines changed: 16 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -25,12 +25,14 @@ class CUDACache {
2525
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_depthDownsampled, other->m_cache[frameFrom].d_depthDownsampled,sizeof(float) * m_width * m_height, cudaMemcpyDeviceToDevice));
2626
//MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_colorDownsampled, other->m_cache[frameFrom].d_colorDownsampled, sizeof(uchar4) * m_width * m_height, cudaMemcpyDeviceToDevice));
2727
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_cameraposDownsampled, other->m_cache[frameFrom].d_cameraposDownsampled,sizeof(float4) * m_width * m_height, cudaMemcpyDeviceToDevice));
28-
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_normalsDownsampled, other->m_cache[frameFrom].d_normalsDownsampled,sizeof(float4) * m_width * m_height, cudaMemcpyDeviceToDevice));
2928

3029
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_intensityDownsampled, other->m_cache[frameFrom].d_intensityDownsampled,sizeof(float) * m_width * m_height, cudaMemcpyDeviceToDevice));
3130
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_intensityDerivsDownsampled, other->m_cache[frameFrom].d_intensityDerivsDownsampled,sizeof(float2) * m_width * m_height, cudaMemcpyDeviceToDevice));
32-
31+
#ifdef CUDACACHE_UCHAR_NORMALS
3332
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_normalsDownsampledUCHAR4, other->m_cache[frameFrom].d_normalsDownsampledUCHAR4,sizeof(float) * m_width * m_height, cudaMemcpyDeviceToDevice));
33+
#else
34+
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[m_currentFrame].d_normalsDownsampled, other->m_cache[frameFrom].d_normalsDownsampled,sizeof(float4) * m_width * m_height, cudaMemcpyDeviceToDevice));
35+
#endif
3436
m_currentFrame++;
3537
}
3638

@@ -69,11 +71,14 @@ class CUDACache {
6971
const CUDACachedFrame& f = m_cache[i];
7072
MLIB_CUDA_SAFE_CALL(cudaMemcpy(depth.getData(), f.d_depthDownsampled,sizeof(float)*depth.getNumPixels(), cudaMemcpyDeviceToHost));
7173
MLIB_CUDA_SAFE_CALL(cudaMemcpy(camPos.getData(), f.d_cameraposDownsampled,sizeof(float4)*camPos.getNumPixels(), cudaMemcpyDeviceToHost));
72-
MLIB_CUDA_SAFE_CALL(cudaMemcpy(normals.getData(), f.d_normalsDownsampled,sizeof(float4)*normals.getNumPixels(), cudaMemcpyDeviceToHost));
7374
//MLIB_CUDA_SAFE_CALL(cudaMemcpy(color.getData(), f.d_colorDownsampled, sizeof(uchar4)*color.getNumPixels(), cudaMemcpyDeviceToHost));
7475
MLIB_CUDA_SAFE_CALL(cudaMemcpy(intensity.getData(), f.d_intensityDownsampled,sizeof(float)*intensity.getNumPixels(), cudaMemcpyDeviceToHost));
7576
MLIB_CUDA_SAFE_CALL(cudaMemcpy(intensityDerivative.getData(), f.d_intensityDerivsDownsampled,sizeof(float2)*intensityDerivative.getNumPixels(), cudaMemcpyDeviceToHost));
77+
#ifdef CUDACACHE_UCHAR_NORMALS
7678
MLIB_CUDA_SAFE_CALL(cudaMemcpy(intensityOrig.getData(), f.d_normalsDownsampledUCHAR4,sizeof(float)*intensityOrig.getNumPixels(), cudaMemcpyDeviceToHost));
79+
#else
80+
MLIB_CUDA_SAFE_CALL(cudaMemcpy(normals.getData(), f.d_normalsDownsampled,sizeof(float4)*normals.getNumPixels(), cudaMemcpyDeviceToHost));
81+
#endif
7782
s << depth;
7883
s << camPos;
7984
s << normals;
@@ -121,11 +126,14 @@ class CUDACache {
121126
s >> intensityOrig;
122127
MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_depthDownsampled, depth.getData(),sizeof(float)*depth.getNumPixels(), cudaMemcpyHostToDevice));
123128
MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_cameraposDownsampled, camPos.getData(),sizeof(float4)*camPos.getNumPixels(), cudaMemcpyHostToDevice));
124-
MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_normalsDownsampled, normals.getData(),sizeof(float4)*normals.getNumPixels(), cudaMemcpyHostToDevice));
125129
//MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_colorDownsampled, color.getData(), sizeof(uchar4)*color.getNumPixels(), cudaMemcpyHostToDevice));
126130
MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_intensityDownsampled, intensity.getData(),sizeof(float)*intensity.getNumPixels(), cudaMemcpyHostToDevice));
127131
MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_intensityDerivsDownsampled, intensityDerivative.getData(),sizeof(float2)*intensityDerivative.getNumPixels(), cudaMemcpyHostToDevice));
132+
#ifdef CUDACACHE_UCHAR_NORMALS
128133
MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_normalsDownsampledUCHAR4, intensityOrig.getData(),sizeof(float)*intensityOrig.getNumPixels(), cudaMemcpyHostToDevice));
134+
#else
135+
MLIB_CUDA_SAFE_CALL(cudaMemcpy(f.d_normalsDownsampled, normals.getData(),sizeof(float4)*normals.getNumPixels(), cudaMemcpyHostToDevice));
136+
#endif
129137
}
130138
s.close();
131139
}
@@ -144,12 +152,14 @@ class CUDACache {
144152
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_depthDownsampled, cachedFrames[i].d_depthDownsampled,sizeof(float) * m_width * m_height, cudaMemcpyDeviceToDevice));
145153
//MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_colorDownsampled, cachedFrames[i].d_colorDownsampled, sizeof(uchar4) * m_width * m_height, cudaMemcpyDeviceToDevice));
146154
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_cameraposDownsampled, cachedFrames[i].d_cameraposDownsampled,sizeof(float4) * m_width * m_height, cudaMemcpyDeviceToDevice));
147-
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_normalsDownsampled, cachedFrames[i].d_normalsDownsampled,sizeof(float4) * m_width * m_height, cudaMemcpyDeviceToDevice));
148155

149156
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_intensityDownsampled, cachedFrames[i].d_intensityDownsampled,sizeof(float) * m_width * m_height, cudaMemcpyDeviceToDevice));
150157
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_intensityDerivsDownsampled, cachedFrames[i].d_intensityDerivsDownsampled,sizeof(float2) * m_width * m_height, cudaMemcpyDeviceToDevice));
151-
158+
#ifdef CUDACACHE_UCHAR_NORMALS
152159
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_normalsDownsampledUCHAR4, cachedFrames[i].d_normalsDownsampledUCHAR4,sizeof(float) * m_width * m_height, cudaMemcpyDeviceToDevice));
160+
#else
161+
MLIB_CUDA_SAFE_CALL(cudaMemcpy(m_cache[i].d_normalsDownsampled, cachedFrames[i].d_normalsDownsampled,sizeof(float4) * m_width * m_height, cudaMemcpyDeviceToDevice));
162+
#endif
153163
}
154164
}
155165

‎FriedLiver/Source/CUDACacheUtil.h‎

Lines changed: 14 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -4,40 +4,48 @@
44

55
#include"mLibCuda.h"
66

7+
#defineCUDACACHE_UCHAR_NORMALS
8+
79
structCUDACachedFrame {
810
voidalloc(unsignedintwidth,unsignedintheight) {
911
MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_depthDownsampled,sizeof(float)*width*height));
1012
//MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_colorDownsampled, sizeof(uchar4) * width * height));
1113
MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_cameraposDownsampled,sizeof(float4)*width*height));
12-
MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_normalsDownsampled,sizeof(float4)*width*height));
1314

1415
MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_intensityDownsampled,sizeof(float)*width*height));
1516
MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_intensityDerivsDownsampled,sizeof(float2)*width*height));
16-
17+
#ifdefCUDACACHE_UCHAR_NORMALS
1718
MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_normalsDownsampledUCHAR4,sizeof(uchar4)*width*height));
19+
#else
20+
MLIB_CUDA_SAFE_CALL(cudaMalloc(&d_normalsDownsampled,sizeof(float4)*width*height));
21+
#endif
1822
}
1923
voidfree() {
2024
MLIB_CUDA_SAFE_FREE(d_depthDownsampled);
2125
//MLIB_CUDA_SAFE_FREE(d_colorDownsampled);
2226
MLIB_CUDA_SAFE_FREE(d_cameraposDownsampled);
23-
MLIB_CUDA_SAFE_FREE(d_normalsDownsampled);
2427

2528
MLIB_CUDA_SAFE_FREE(d_intensityDownsampled);
2629
MLIB_CUDA_SAFE_FREE(d_intensityDerivsDownsampled);
27-
30+
#ifdefCUDACACHE_UCHAR_NORMALS
2831
MLIB_CUDA_SAFE_FREE(d_normalsDownsampledUCHAR4);
32+
#else
33+
MLIB_CUDA_SAFE_FREE(d_normalsDownsampled);
34+
#endif
2935
}
3036

3137
float*d_depthDownsampled;
3238
//uchar4* d_colorDownsampled;
3339
float4*d_cameraposDownsampled;
34-
float4*d_normalsDownsampled;
3540

3641
//for dense color term
3742
float*d_intensityDownsampled;//this could be packed with intensityDerivaties to a float4 dunno about the read there
3843
float2*d_intensityDerivsDownsampled;//TODO could have energy over intensity gradient instead of intensity
39-
44+
#ifdefCUDACACHE_UCHAR_NORMALS
4045
uchar4*d_normalsDownsampledUCHAR4;
46+
#else
47+
float4*d_normalsDownsampled;
48+
#endif
4149
};
4250

4351
#endif//CUDA_CACHE_UTIL

‎FriedLiver/Source/OnlineBundler.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -88,7 +88,7 @@ class OnlineBundler {
8888
std::mutexmutex_completeTrajectory;
8989
float4x4*d_completeTrajectory;
9090
float4x4*d_localTrajectories;
91-
std::vector<std::vector<int>> m_localTrajectoriesValid;//TODO CHECK IF NEEDED
91+
std::vector<std::vector<int>> m_localTrajectoriesValid;
9292

9393
float4x4*d_siftTrajectory;// frame-to-frame sift tracking for all frames in sequence
9494
//************************************

‎FriedLiver/Source/SBA.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -103,7 +103,7 @@ bool SBA::align(SIFTImageManager* siftManager, const CUDACache* cudaCache, float
103103
if (!isScanDoneOpt &&GlobalBundlingState::get().s_enableGlobalTimings) {cudaDeviceSynchronize(); s_timer.start(); }
104104

105105
unsignedint numImages = siftManager->getNumImages();
106-
//if (isStart) siftManager->updateGPUValidImages(); //TODO CHECK
106+
//if (isStart) siftManager->updateGPUValidImages(); //should be in sync already
107107
constint* d_validImages = siftManager->getValidImagesGPU();
108108
convertMatricesToPosesCU(d_transforms, numImages, d_xRot, d_xTrans, d_validImages);
109109

‎FriedLiver/Source/SiftGPU/SIFTImageManager.cpp‎

Lines changed: 1 addition & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -553,12 +553,6 @@ unsigned int SIFTImageManager::filterFrames(unsigned int curFrame, unsigned int
553553
std::vector<unsignedint>currNumFilteredMatchesPerImagePair(numFrames - startFrame);
554554
MLIB_CUDA_SAFE_CALL(cudaMemcpy(currNumFilteredMatchesPerImagePair.data(), d_currNumFilteredMatchesPerImagePair + startFrame,sizeof(unsignedint) * currNumFilteredMatchesPerImagePair.size(), cudaMemcpyDeviceToHost));
555555

556-
//for (unsigned int i = startFrame; i < numFrames; i++) { // previous frames
557-
//if (m_validImages[i] != 0 && currNumFilteredMatchesPerImagePair[i-startFrame] > 0 && i != curFrame) {
558-
//connected = 1;
559-
//break;
560-
//}
561-
//}
562556
for (int i = (int)numFrames-1; i >= (int)startFrame; i--) {// previous frames (reverse direction)
563557
if (m_validImages[i] !=0 && currNumFilteredMatchesPerImagePair[i - startFrame] >0 && i != curFrame) {
564558
connected =1;
@@ -571,7 +565,7 @@ unsigned int SIFTImageManager::filterFrames(unsigned int curFrame, unsigned int
571565
//std::cout << "frame " << curFrame << " not connected to previous!" << std::endl;
572566

573567
m_validImages[curFrame] = connected;
574-
MLIB_CUDA_SAFE_CALL(cudaMemcpy(d_validImages + curFrame, &m_validImages[curFrame],sizeof(int), cudaMemcpyHostToDevice));//TODO CHECK
568+
MLIB_CUDA_SAFE_CALL(cudaMemcpy(d_validImages + curFrame, &m_validImages[curFrame],sizeof(int), cudaMemcpyHostToDevice));
575569
return lastMatchedFrame;
576570
}
577571

‎FriedLiver/Source/SiftGPU/SIFTImageManager.cu‎

Lines changed: 36 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -414,16 +414,30 @@ void SIFTImageManager::FilterMatchesBySurfaceAreaCU(unsigned int curFrame, unsig
414414

415415
#defineFILTER_DENSE_VERIFY_THREAD_SPLIT32
416416

417+
#ifdef CUDACACHE_UCHAR_NORMALS
418+
__device__float3computeProjError(unsignedint idx,unsignedint imageWidth,unsignedint imageHeight,
419+
float distThresh,float normalThresh,float colorThresh,const float4x4& transform,const float4x4& intrinsics,
420+
constfloat* d_inputDepth,constfloat4* d_inputCamPos,constuchar4* d_inputNormal,constfloat* d_inputColor,
421+
constfloat* d_modelDepth,constfloat4* d_modelCamPos,constuchar4* d_modelNormal,constfloat* d_modelColor,
422+
float sensorDepthMin,float sensorDepthMax)
423+
#else
417424
__device__float3computeProjError(unsignedint idx,unsignedint imageWidth,unsignedint imageHeight,
418425
float distThresh,float normalThresh,float colorThresh,const float4x4& transform,const float4x4& intrinsics,
419426
constfloat* d_inputDepth,constfloat4* d_inputCamPos,constfloat4* d_inputNormal,constfloat* d_inputColor,
420427
constfloat* d_modelDepth,constfloat4* d_modelCamPos,constfloat4* d_modelNormal,constfloat* d_modelColor,
421428
float sensorDepthMin,float sensorDepthMax)
429+
#endif
422430
{
423431
float3 out =make_float3(0.0f);
424432

425433
float4 pInput = d_inputCamPos[idx];// point
434+
#ifdef CUDACACHE_UCHAR_NORMALS
435+
float4 nInput =make_float4(MINF);
436+
uchar4 nInputU4 = d_inputNormal[idx];
437+
if (*(int*)(&nInputU4) !=0) nInput =make_float4(make_float3(nInputU4.x, nInputU4.y, nInputU4.z) /255.0f *2.0f -1.0f,0.0f);// vector
438+
#else
426439
float4 nInput = d_inputNormal[idx]; nInput.w =0.0f;// vector
440+
#endif
427441
float dInput = d_inputDepth[idx];
428442
//float cInput = d_inputColor[idx];
429443

@@ -436,9 +450,14 @@ __device__ float3 computeProjError(unsigned int idx, unsigned int imageWidth, un
436450

437451
if (screenPos.x >=0 && screenPos.y >=0 && screenPos.x < (int)imageWidth && screenPos.y < (int)imageHeight) {
438452
float4 pTarget = d_modelCamPos[screenPos.y * imageWidth + screenPos.x];//getBestCorrespondence1x1
439-
float4 nTarget = d_modelNormal[screenPos.y * imageWidth + screenPos.x];
440453
//float cTarget = d_modelColor[screenPos.y * imageWidth + screenPos.x];
441-
454+
#ifdef CUDACACHE_UCHAR_NORMALS
455+
float4 nTarget =make_float4(MINF);
456+
uchar4 nTargetU4 = d_modelNormal[idx];
457+
if (*(int*)(&nTargetU4) !=0) nTarget =make_float4(make_float3(nTargetU4.x, nTargetU4.y, nTargetU4.z) /255.0f *2.0f -1.0f,0.0f);// vector
458+
#else
459+
float4 nTarget = d_modelNormal[screenPos.y * imageWidth + screenPos.x];
460+
#endif
442461
if (pTarget.x != MINF && nTarget.x != MINF) {
443462
float d =length(pTransInput - pTarget);
444463
float dNormal =dot(make_float3(nTransInput.x, nTransInput.y, nTransInput.z),make_float3(nTarget.x, nTarget.y, nTarget.z));// should be able to do dot(nTransInput, nTarget)
@@ -483,14 +502,18 @@ void __global__ FilterMatchesByDenseVerifyCU_Kernel(unsigned int curImageIdx, un
483502

484503
constfloat* d_inputDepth = d_cachedFrames[imagePairIdx].d_depthDownsampled;
485504
constfloat4* d_inputCamPos = d_cachedFrames[imagePairIdx].d_cameraposDownsampled;
486-
constfloat4* d_inputNormal = d_cachedFrames[imagePairIdx].d_normalsDownsampled;
487505
constfloat* d_inputColor = d_cachedFrames[imagePairIdx].d_intensityDownsampled;
488506

489507
constfloat* d_modelDepth = d_cachedFrames[curImageIdx].d_depthDownsampled;
490508
constfloat4* d_modelCamPos = d_cachedFrames[curImageIdx].d_cameraposDownsampled;
491-
constfloat4* d_modelNormal = d_cachedFrames[curImageIdx].d_normalsDownsampled;
492509
constfloat* d_modelColor = d_cachedFrames[curImageIdx].d_intensityDownsampled;
493-
510+
#ifdef CUDACACHE_UCHAR_NORMALS
511+
constuchar4* d_inputNormal = d_cachedFrames[imagePairIdx].d_normalsDownsampledUCHAR4;
512+
constuchar4* d_modelNormal = d_cachedFrames[curImageIdx].d_normalsDownsampledUCHAR4;
513+
#else
514+
constfloat4* d_inputNormal = d_cachedFrames[imagePairIdx].d_normalsDownsampled;
515+
constfloat4* d_modelNormal = d_cachedFrames[curImageIdx].d_normalsDownsampled;
516+
#endif
494517
const float4x4 transform = d_currFilteredTransforms[imagePairIdx];
495518

496519

@@ -1023,14 +1046,20 @@ void __global__ VerifyTrajectoryCU_Kernel(unsigned int numImages, int* d_validIm
10231046

10241047
constfloat* d_inputDepth = d_cachedFrames[img0].d_depthDownsampled;
10251048
constfloat4* d_inputCamPos = d_cachedFrames[img0].d_cameraposDownsampled;
1026-
constfloat4* d_inputNormal = d_cachedFrames[img0].d_normalsDownsampled;
10271049
constfloat* d_inputColor = d_cachedFrames[img0].d_intensityDownsampled;
10281050

10291051
constfloat* d_modelDepth = d_cachedFrames[img1].d_depthDownsampled;
10301052
constfloat4* d_modelCamPos = d_cachedFrames[img1].d_cameraposDownsampled;
1031-
constfloat4* d_modelNormal = d_cachedFrames[img1].d_normalsDownsampled;
10321053
constfloat* d_modelColor = d_cachedFrames[img1].d_intensityDownsampled;
10331054

1055+
#ifdef CUDACACHE_UCHAR_NORMALS
1056+
constuchar4* d_inputNormal = d_cachedFrames[img0].d_normalsDownsampledUCHAR4;
1057+
constuchar4* d_modelNormal = d_cachedFrames[img1].d_normalsDownsampledUCHAR4;
1058+
#else
1059+
constfloat4* d_inputNormal = d_cachedFrames[img0].d_normalsDownsampled;
1060+
constfloat4* d_modelNormal = d_cachedFrames[img1].d_normalsDownsampled;
1061+
#endif
1062+
10341063
const float4x4 transform = d_trajectory[img1].getInverse() * d_trajectory[img0];
10351064

10361065
float local_sumResidual =0.0f;

‎FriedLiver/Source/SiftGPU/SIFTMatchFilter.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -186,6 +186,9 @@ float2 SIFTMatchFilter::computeSurfaceArea(const SIFTKeyPoint* keys, const uint2
186186

187187
voidSIFTMatchFilter::filterByDenseVerify(SIFTImageManager* siftManager,const std::vector<CUDACachedFrame>& cachedFrames,const float4x4& depthIntrinsics,float depthMin,float depthMax)
188188
{
189+
#ifdef CUDACACHE_UCHAR_NORMALS
190+
MLIB_EXCEPTION("need to update to uchar4 normals");
191+
#else
189192
constunsignedint numImages = siftManager->getNumImages();
190193
if (numImages <=1)return;
191194

@@ -235,6 +238,7 @@ void SIFTMatchFilter::filterByDenseVerify(SIFTImageManager* siftManager, const s
235238
cutilSafeCall(cudaMemcpy(siftManager->d_currNumFilteredMatchesPerImagePair + i, &numMatches,sizeof(unsignedint), cudaMemcpyHostToDevice));
236239
}
237240
}
241+
#endif
238242
}
239243

240244
boolSIFTMatchFilter::filterImagePairByDenseVerify(constfloat* inputDepth,const float4* inputCamPos,const float4* inputNormals,constfloat* inputColor,
@@ -713,6 +717,9 @@ void SIFTMatchFilter::filterKeyPointMatchesDEBUG(unsigned int curFrame, SIFTImag
713717
voidSIFTMatchFilter::visualizeProjError(SIFTImageManager* siftManager,const vec2ui& imageIndices,const std::vector<CUDACachedFrame>& cachedFrames,
714718
const float4x4& depthIntrinsics,const float4x4& transformCurToPrv,float depthMin,float depthMax)
715719
{
720+
#ifdef CUDACACHE_UCHAR_NORMALS
721+
throwMLIB_EXCEPTION("need to update to uchar4 normals");
722+
#else
716723
constunsignedint numImages = siftManager->getNumImages();
717724
MLIB_ASSERT(imageIndices.x < numImages && imageIndices.y < numImages);
718725

@@ -788,6 +795,7 @@ void SIFTMatchFilter::visualizeProjError(SIFTImageManager* siftManager, const ve
788795
if (projErrors.x == -std::numeric_limits<float>::infinity() || (projErrors.x > verifyOptErrThresh) || (projErrors.y < verifyOptCorrThresh)) {// tracking lost or bad match
789796
valid =false;// invalid
790797
}
798+
#endif
791799
}
792800

793801
voidSIFTMatchFilter::computeCorrespondencesDEBUG(unsignedint width,unsignedint height,constfloat* inputDepth,const float4* input,const float4* inputNormals,constfloat* inputColor,

0 commit comments

Comments
 (0)
[8]ページ先頭
©2009-2025 Movatter.jp