From d5e5067acd4e7a3a98f3c96e894c591d85fa336a Mon Sep 17 00:00:00 2001 From: Minyan Zhong Date: Mon, 11 Jan 2021 02:15:22 -0800 Subject: [PATCH] update --- contrib/PyCuAmpcor/src/cuAmpcorChunk.cu | 2 +- contrib/PyCuAmpcor/src/cuAmpcorUtil.h | 2 +- contrib/PyCuAmpcor/src/cuEstimateStats.cu | 20 ++++++++++---------- 3 files changed, 12 insertions(+), 12 deletions(-) diff --git a/contrib/PyCuAmpcor/src/cuAmpcorChunk.cu b/contrib/PyCuAmpcor/src/cuAmpcorChunk.cu index 4ec7f38..286b099 100644 --- a/contrib/PyCuAmpcor/src/cuAmpcorChunk.cu +++ b/contrib/PyCuAmpcor/src/cuAmpcorChunk.cu @@ -66,7 +66,7 @@ void cuAmpcorChunk::run(int idxDown_, int idxAcross_) cuArraysMaxloc2D(r_corrBatchRaw, offsetInit, r_maxval, stream); // estimate variance - cuEstimateVariance(r_referenceBatchRaw->size, r_corrBatchRaw, offsetInit, r_maxval, r_covValue, stream); + cuEstimateVariance(r_corrBatchRaw, offsetInit, r_maxval, r_referenceBatchRaw->size, r_covValue, stream); // estimate SNR // step1: extraction of correlation surface around the peak diff --git a/contrib/PyCuAmpcor/src/cuAmpcorUtil.h b/contrib/PyCuAmpcor/src/cuAmpcorUtil.h index bdfbaa1..55961f8 100644 --- a/contrib/PyCuAmpcor/src/cuAmpcorUtil.h +++ b/contrib/PyCuAmpcor/src/cuAmpcorUtil.h @@ -91,7 +91,7 @@ void cuArraysSumCorr(cuArrays *images, cuArrays *imagesValid, cuArra void cuEstimateSnr(cuArrays *corrSum, cuArrays *corrValidCount, cuArrays *maxval, cuArrays *snrValue, cudaStream_t stream); // implemented in cuEstimateStats.cu -void cuEstimateVariance(int winSize, cuArrays *corrBatchRaw, cuArrays *maxloc, cuArrays *maxval, cuArrays *covValue, cudaStream_t stream); +void cuEstimateVariance(cuArrays *corrBatchRaw, cuArrays *maxloc, cuArrays *maxval, int templateSize, cuArrays *covValue, cudaStream_t stream); #endif diff --git a/contrib/PyCuAmpcor/src/cuEstimateStats.cu b/contrib/PyCuAmpcor/src/cuEstimateStats.cu index 0a24420..7d31673 100644 --- a/contrib/PyCuAmpcor/src/cuEstimateStats.cu +++ b/contrib/PyCuAmpcor/src/cuEstimateStats.cu @@ -46,8 +46,8 @@ void cuEstimateSnr(cuArrays *corrSum, cuArrays *corrValidCount, cuAr } // cuda kernel for cuEstimateVariance -__global__ void cudaKernel_estimateVar(const int winSize, const float* corrBatchRaw, const int NX, const int NY, - const int2* maxloc, const float* maxval, float3* covValue, const int size) +__global__ void cudaKernel_estimateVar(const float* corrBatchRaw, const int NX, const int NY, const int2* maxloc, + const float* maxval, const int templateSize, float3* covValue, const int size) { // Find image id. @@ -63,7 +63,7 @@ __global__ void cudaKernel_estimateVar(const int winSize, const float* corrBatch // Check if maxval is on the margin. if (px-1 < 0 || py-1 <0 || px + 1 >=NX || py+1 >=NY) { - covValue[idxImage] = make_float3(99.0, 99.0, 99.0); + covValue[idxImage] = make_float3(99.0, 99.0, 0.0); } else { @@ -84,13 +84,13 @@ __global__ void cudaKernel_estimateVar(const int winSize, const float* corrBatch float n2 = fmaxf(1 - peak, 0.0); - dxx = dxx * winSize; - dyy = dyy * winSize; - dxy = dxy * winSize; + dxx = dxx * templateSize; + dyy = dyy * templateSize; + dxy = dxy * templateSize; float n4 = n2*n2; n2 = n2 * 2; - n4 = n4 * 0.5 * winSize; + n4 = n4 * 0.5 * templateSize; float u = dxy * dxy - dxx * dyy; float u2 = u*u; @@ -111,19 +111,19 @@ __global__ void cudaKernel_estimateVar(const int winSize, const float* corrBatch /** * Estimate the variance of the correlation surface - * @param[in] winSize size of reference chip + * @param[in] templateSize size of reference chip * @param[in] corrBatchRaw correlation surface * @param[in] maxloc maximum location * @param[in] maxval maximum value * @param[out] covValue variance value * @param[in] stream cuda stream */ -void cuEstimateVariance(int winSize, cuArrays *corrBatchRaw, cuArrays *maxloc, cuArrays *maxval, cuArrays *covValue, cudaStream_t stream) +void cuEstimateVariance(cuArrays *corrBatchRaw, cuArrays *maxloc, cuArrays *maxval, int templateSize, cuArrays *covValue, cudaStream_t stream) { int size = corrBatchRaw->count; // One dimensional launching parameters to loop over every correlation surface. cudaKernel_estimateVar<<< IDIVUP(size, NTHREADS), NTHREADS, 0, stream>>> - (winSize, corrBatchRaw->devData, corrBatchRaw->height, corrBatchRaw->width, maxloc->devData, maxval->devData, covValue->devData, size); + (corrBatchRaw->devData, corrBatchRaw->height, corrBatchRaw->width, maxloc->devData, maxval->devData, templateSize, covValue->devData, size); getLastCudaError("cudaKernel_estimateVar error\n"); } //end of file