Skip to content

Commit 269759b

Browse files
Liang YuGitHub Enterprise
authored andcommitted
Fix resample chip phase computation (#931)
* fix hardcoded indexing in chip phase calculation
1 parent 0d82a80 commit 269759b

File tree

2 files changed

+25
-26
lines changed

2 files changed

+25
-26
lines changed

cxx/isce3/cuda/image/gpuResampSlc.cu

Lines changed: 23 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -53,11 +53,11 @@ void transformTile(const thrust::complex<float> *tile,
5353
double refRangePixelSpacing,
5454
double refWavelength,
5555
int chipSize,
56-
int rowOffset,
56+
int rowOffset,
5757
int rowStart) {
5858

5959
int iTileOut = blockDim.x * blockIdx.x + threadIdx.x;
60-
int iChip = iTileOut * chipSize * chipSize;
60+
int iChip = iTileOut * chipSize * chipSize;
6161
int chipHalf = chipSize/2;
6262

6363
if (iTileOut < outWidth*outLength) {
@@ -74,7 +74,7 @@ void transformTile(const thrust::complex<float> *tile,
7474
const int intRg = __float2int_rd(j + rgOff);
7575
const double fracAz = i + azOff - intAz + rowStart;
7676
const double fracRg = j + rgOff - intRg;
77-
77+
7878
// Check bounds again. Use rowOffset to account tiles where tile.rowStart != tile.firstRowImage
7979
bool intAzInBounds = !((intAz+rowOffset < chipHalf) || (intAz >= (inLength - chipHalf)));
8080
bool intRgInBounds = !((intRg < chipHalf) || (intRg >= (inWidth - chipHalf)));
@@ -86,29 +86,29 @@ void transformTile(const thrust::complex<float> *tile,
8686

8787
// Doppler to be added back. Simultaneously evaluate carrier that needs to
8888
// be added back after interpolation
89-
double phase = (dop * fracAz)
90-
+ rgCarrier.eval(i + azOff, j + rgOff)
89+
double phase = (dop * fracAz)
90+
+ rgCarrier.eval(i + azOff, j + rgOff)
9191
+ azCarrier.eval(i + azOff, j + rgOff);
9292

9393
// Flatten the carrier phase if requested
9494
if (flatten) {
95-
phase += ((4. * (M_PI / wavelength)) *
96-
((startingRange - refStartingRange)
97-
+ (j * (rangePixelSpacing - refRangePixelSpacing))
98-
+ (rgOff * rangePixelSpacing))) + ((4.0 * M_PI
99-
* (refStartingRange + (j * refRangePixelSpacing)))
95+
phase += ((4. * (M_PI / wavelength)) *
96+
((startingRange - refStartingRange)
97+
+ (j * (rangePixelSpacing - refRangePixelSpacing))
98+
+ (rgOff * rangePixelSpacing))) + ((4.0 * M_PI
99+
* (refStartingRange + (j * refRangePixelSpacing)))
100100
* ((1.0 / refWavelength) - (1.0 / wavelength)));
101101
}
102-
102+
103103
// Modulate by 2*PI
104104
phase = fmod(phase, 2.0*M_PI);
105-
105+
106106
// Read data chip without the carrier phases
107107
for (int ii = 0; ii < chipSize; ++ii) {
108108
// Row to read from
109109
const int chipRow = intAz + ii - chipHalf + rowOffset - rowStart;
110110
// Carrier phase
111-
const double phase = dop * (ii - 4.0);
111+
const double phase = dop * (ii - chipHalf);
112112
const thrust::complex<float> cval(cos(phase), -sin(phase));
113113
// Set the data values after removing doppler in azimuth
114114
for (int jj = 0; jj < chipSize; ++jj) {
@@ -142,7 +142,7 @@ gpuTransformTile(isce3::image::Tile<std::complex<float>> & tile,
142142
isce3::cuda::core::gpuSinc2dInterpolator<thrust::complex<float>> interp,
143143
int inWidth, int inLength, double startingRange, double rangePixelSpacing,
144144
double prf, double wavelength, double refStartingRange,
145-
double refRangePixelSpacing, double refWavelength,
145+
double refRangePixelSpacing, double refWavelength,
146146
bool flatten, int chipSize) {
147147

148148
// Cache geometry values
@@ -186,14 +186,14 @@ gpuTransformTile(isce3::image::Tile<std::complex<float>> & tile,
186186
dim3 grid((nOutPixels+(THRD_PER_BLOCK-1))/THRD_PER_BLOCK);
187187

188188
// global call to transform
189-
transformTile<<<grid, block>>>(d_tile,
189+
transformTile<<<grid, block>>>(d_tile,
190190
d_chip,
191-
d_imgOut,
192-
d_rgOffTile,
193-
d_azOffTile,
194-
d_rgCarrier,
195-
d_azCarrier,
196-
d_dopplerLUT,
191+
d_imgOut,
192+
d_rgOffTile,
193+
d_azOffTile,
194+
d_rgCarrier,
195+
d_azCarrier,
196+
d_dopplerLUT,
197197
interp,
198198
flatten,
199199
outWidth,
@@ -206,7 +206,7 @@ gpuTransformTile(isce3::image::Tile<std::complex<float>> & tile,
206206
wavelength,
207207
refStartingRange,
208208
refRangePixelSpacing,
209-
refWavelength,
209+
refWavelength,
210210
chipSize,
211211
tile.rowStart()-tile.firstImageRow(),// needed to keep az in bounds in subtiles
212212
tile.rowStart()); // needed to match az components on CPU
@@ -224,7 +224,7 @@ gpuTransformTile(isce3::image::Tile<std::complex<float>> & tile,
224224
checkCudaErrors(cudaFree(d_imgOut));
225225
checkCudaErrors(cudaFree(d_azOffTile));
226226
checkCudaErrors(cudaFree(d_rgOffTile));
227-
227+
228228
// Write block of data
229229
outputSlc.setBlock(imgOut, 0, tile.rowStart(), outWidth, outLength);
230230
}

cxx/isce3/image/ResampSlc.cpp

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -322,7 +322,7 @@ void ResampSlc::_transformTile(Tile_t& tile, Raster& outputSlc,
322322
const int chipRow =
323323
intAz - tile.firstImageRow() + ii - chipHalf;
324324
// Carrier phase
325-
const double chipPhase = dop * (ii - 4.0);
325+
const double chipPhase = dop * (ii - chipHalf);
326326
const std::complex<float> cval(
327327
std::cos(chipPhase), -std::sin(chipPhase));
328328
// Set the data values after removing doppler in azimuth
@@ -335,8 +335,7 @@ void ResampSlc::_transformTile(Tile_t& tile, Raster& outputSlc,
335335

336336
// Interpolate chip
337337
const std::complex<float> cval = _interp->interpolate(
338-
isce3::core::SINC_HALF + fracRg,
339-
isce3::core::SINC_HALF + fracAz, chip);
338+
chipHalf + fracRg, chipHalf + fracAz, chip);
340339

341340
// Add doppler to interpolated value and save
342341
imgOut[tileLine * outWidth + j] =

0 commit comments

Comments
 (0)