Skip to content

Commit 76660df

Browse files
Liang YuGitHub Enterprise
authored andcommitted
CUDA resamp LUT2d undo (#992)
* CUDA resample SLC doppler temporarily reverted back to LUT1d
1 parent 36ab7f4 commit 76660df

File tree

3 files changed

+12
-11
lines changed

3 files changed

+12
-11
lines changed

cxx/isce3/cuda/image/ResampSlc.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77

88
// isce3::core
99
#include <isce3/core/Constants.h>
10-
#include <isce3/core/LUT2d.h>
10+
#include <isce3/core/LUT1d.h>
1111

1212
#include <isce3/image/Tile.h>
1313

@@ -75,6 +75,7 @@ resamp(isce3::io::Raster & inputSlc, isce3::io::Raster & outputSlc,
7575
auto timerStart = std::chrono::steady_clock::now();
7676

7777
// For each full tile of _linesPerTile lines...
78+
const isce3::core::LUT1d<double> dopplerLUT1d = isce3::core::avgLUT2dToLUT1d<double>(_dopplerLUT);
7879
for (int tileCount = 0; tileCount < nTiles; tileCount++) {
7980

8081
// Make a tile for representing input SLC data
@@ -100,7 +101,8 @@ resamp(isce3::io::Raster & inputSlc, isce3::io::Raster & outputSlc,
100101
// Perform interpolation
101102
std::cout << "Interpolating tile " << tileCount << std::endl;
102103
gpuTransformTile(tile, outputSlc, rgOffTile, azOffTile, _rgCarrier, _azCarrier,
103-
_dopplerLUT, interp, inWidth, inLength, this->startingRange(),
104+
dopplerLUT1d, interp,
105+
inWidth, inLength, this->startingRange(),
104106
this->rangePixelSpacing(), this->sensingStart(), this->prf(),
105107
this->wavelength(), this->refStartingRange(),
106108
this->refRangePixelSpacing(), this->refWavelength(), flatten,

cxx/isce3/cuda/image/gpuResampSlc.cu

Lines changed: 7 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111

1212
// isce3::cuda::core
1313
#include <isce3/cuda/core/gpuPoly2d.h>
14-
#include <isce3/cuda/core/gpuLUT2d.h>
14+
#include <isce3/cuda/core/gpuLUT1d.h>
1515
#include <isce3/cuda/core/gpuInterpolator.h>
1616

1717
#include <isce3/cuda/except/Error.h>
@@ -20,7 +20,7 @@
2020

2121
using isce3::cuda::core::gpuPoly2d;
2222
using isce3::cuda::core::gpuInterpolator;
23-
using isce3::cuda::core::gpuLUT2d;
23+
using isce3::cuda::core::gpuLUT1d;
2424
using isce3::cuda::core::gpuSinc2dInterpolator;
2525

2626
#define THRD_PER_BLOCK 512// Number of threads per block (should always %32==0)
@@ -33,7 +33,7 @@ void transformTile(const thrust::complex<float> *tile,
3333
const float *azOffTile,
3434
const gpuPoly2d rgCarrier,
3535
const gpuPoly2d azCarrier,
36-
const gpuLUT2d<double> dopplerLUT,
36+
const gpuLUT1d<double> dopplerLUT,
3737
gpuSinc2dInterpolator<thrust::complex<float>> interp,
3838
bool flatten,
3939
int outWidth,
@@ -85,12 +85,11 @@ void transformTile(const thrust::complex<float> *tile,
8585

8686
// Skip computations if indices out of bound or az/rng not in doppler
8787
// Output previously filled with designated invalid values
88-
if (intAzOutOfBounds || intRgOutOfBounds
89-
|| !dopplerLUT.contains(az, rng))
88+
if (intAzOutOfBounds || intRgOutOfBounds)
9089
return;
9190

9291
// evaluate Doppler polynomial
93-
const double dop = dopplerLUT.eval(az, rng) * 2 * M_PI / prf;
92+
const double dop = dopplerLUT.eval(rng) * 2 * M_PI / prf;
9493

9594
// Doppler to be added back. Simultaneously evaluate carrier that needs to
9695
// be added back after interpolation
@@ -147,7 +146,7 @@ gpuTransformTile(isce3::image::Tile<std::complex<float>> & tile,
147146
isce3::image::Tile<float> & azOffTile,
148147
const isce3::core::Poly2d & rgCarrier,
149148
const isce3::core::Poly2d & azCarrier,
150-
const isce3::core::LUT2d<double> & dopplerLUT,
149+
const isce3::core::LUT1d<double> & dopplerLUT,
151150
isce3::cuda::core::gpuSinc2dInterpolator<thrust::complex<float>> interp,
152151
int inWidth, int inLength, double startingRange, double rangePixelSpacing,
153152
double sensingStart, double prf, double wavelength, double refStartingRange,
@@ -171,7 +170,7 @@ gpuTransformTile(isce3::image::Tile<std::complex<float>> & tile,
171170
float *d_rgOffTile, *d_azOffTile;
172171
gpuPoly2d d_rgCarrier(rgCarrier);
173172
gpuPoly2d d_azCarrier(azCarrier);
174-
gpuLUT2d<double> d_dopplerLUT(dopplerLUT);
173+
gpuLUT1d<double> d_dopplerLUT(dopplerLUT);
175174

176175
// determine sizes
177176
size_t nInPixels = size_t(tile.length()) * tile.width();

cxx/isce3/cuda/image/gpuResampSlc.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ namespace isce3 {
2929
isce3::image::Tile<float> & azOffTile,
3030
const isce3::core::Poly2d & rgCarrier,
3131
const isce3::core::Poly2d & azCarrier,
32-
const isce3::core::LUT2d<double> & dopplerLUT,
32+
const isce3::core::LUT1d<double> & dopplerLUT,
3333
isce3::cuda::core::gpuSinc2dInterpolator<thrust::complex<float>> interp,
3434
int inWidth, int inLength, double startingRange, double rangePixelSpacing,
3535
double sensingStart, double prf, double wavelength, double refStartingRange,

0 commit comments

Comments
 (0)