Skip to content

Commit 3f62709

Browse files
Virginia BrancatoLiang Yu
authored andcommitted
Integrate SNAPHU in InSAR workflow (#968)
* Move correlation_thresh_increments to common ICU, Phass parameters * Add snaphu parameters in InSAR schema * Re-organize unwrap.py to include snaphu unwrapper * Integrate SNAPHU in InSAR workflow * Correct rg bandwidth, conn components to uint32, fix stats * Add functionality to clip nan from correlation data * add support for unsigned int rasters to CUDA geocode * Fix computation azimuth resolution. Log ICU and Phass attributes * Correct typo in clean cfg function * Rename some functionalities for clarity * Expose and include MCF initializer * add raster data type and interpolation method check in CUDA geocode * Remove default value comments for SNAPHU in InSAR schema * Remove correlation and phase std model from schema and unwrap.py * Fix identation doc strings * Change Snaphu schema params name to match Python wrapper * Use ternary if conditions and fix params in schema * Resolve param name diff between Python interface and InSAR schema * limit int type interp to nearest neighbor * Move unwrap default values to InSAR default * Fix pep8 notation mistakes * remove erroneously introduced file * fix incorrect init logic * added int type and interp type check whitespace cleanup * Include SNAPHU in Rosamond insar test case * Remove function cleaning NaNs * Correct bug on band from where to read data * Move unwrap-related comments from schema to InSAR default runconfig * Correct get function for cfg info extraction * Fix bug of connected component type Co-authored-by: Liang Yu <[email protected]>
1 parent 5176895 commit 3f62709

File tree

15 files changed

+797
-244
lines changed

15 files changed

+797
-244
lines changed

cxx/isce3/cuda/core/InterpolatorHandle.cu

Lines changed: 25 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -16,28 +16,23 @@ namespace isce3::cuda::core {
1616

1717
template<class T>
1818
__global__ void init_interp(DeviceInterp<T>** interp,
19-
isce3::core::dataInterpMethod interp_method, bool* unsupported_interp)
19+
isce3::core::dataInterpMethod interp_method)
2020
{
2121
if (threadIdx.x == 0 && blockIdx.x == 0) {
2222
// Choose interpolator
23-
switch(interp_method) {
24-
case isce3::core::BILINEAR_METHOD:
25-
(*interp) = new isce3::cuda::core::gpuBilinearInterpolator<T>();
26-
break;
27-
case isce3::core::BICUBIC_METHOD:
28-
(*interp) = new isce3::cuda::core::gpuBicubicInterpolator<T>();
29-
break;
30-
case isce3::core::BIQUINTIC_METHOD:
31-
size_t order = 6;
32-
(*interp) = new isce3::cuda::core::gpuSpline2dInterpolator<T>(order);
33-
break;
34-
case isce3::core::NEAREST_METHOD:
35-
(*interp) = new isce3::cuda::core::gpuNearestNeighborInterpolator<T>();
36-
break;
37-
default:
38-
*unsupported_interp = true;
39-
break;
23+
if (interp_method == isce3::core::BILINEAR_METHOD)
24+
(*interp) = new isce3::cuda::core::gpuBilinearInterpolator<T>();
25+
26+
if (interp_method == isce3::core::BICUBIC_METHOD)
27+
(*interp) = new isce3::cuda::core::gpuBicubicInterpolator<T>();
28+
29+
if (interp_method == isce3::core::BIQUINTIC_METHOD) {
30+
size_t order = 6;
31+
(*interp) = new isce3::cuda::core::gpuSpline2dInterpolator<T>(order);
4032
}
33+
34+
if (interp_method == isce3::core::NEAREST_METHOD)
35+
(*interp) = new isce3::cuda::core::gpuNearestNeighborInterpolator<T>();
4136
}
4237
}
4338

@@ -55,21 +50,23 @@ InterpolatorHandle<T>::InterpolatorHandle(
5550
{
5651
checkCudaErrors(cudaMalloc(&_interp, sizeof(DeviceInterp<T>**)));
5752

58-
thrust::device_vector<bool> d_unsupported_interp(1, false);
59-
init_interp<<<1, 1>>>(
60-
_interp, interp_method, d_unsupported_interp.data().get());
61-
checkCudaErrors(cudaPeekAtLastError());
62-
checkCudaErrors(cudaDeviceSynchronize());
63-
64-
bool unsupported_interp = d_unsupported_interp[0];
65-
if (unsupported_interp) {
53+
if (interp_method != isce3::core::BILINEAR_METHOD
54+
&& interp_method != isce3::core::BICUBIC_METHOD
55+
&& interp_method != isce3::core::BIQUINTIC_METHOD
56+
&& interp_method != isce3::core::NEAREST_METHOD)
57+
{
6658
pyre::journal::error_t error(
67-
"isce.cuda.core.InterpolatorHandle.InterpolatorHandle");
59+
"isce3.cuda.core.InterpolatorHandle.InterpolatorHandle");
6860
error << "Unsupported interpolator method provided."
6961
<< pyre::journal::endl;
7062
throw isce3::except::InvalidArgument(
7163
ISCE_SRCINFO(), "Unsupported interpolator method provided.");
7264
}
65+
66+
thrust::device_vector<bool> d_unsupported_interp(1, false);
67+
init_interp<<<1, 1>>>(_interp, interp_method);
68+
checkCudaErrors(cudaPeekAtLastError());
69+
checkCudaErrors(cudaDeviceSynchronize());
7370
}
7471

7572
template<class T>
@@ -86,4 +83,5 @@ template class InterpolatorHandle<thrust::complex<float>>;
8683
template class InterpolatorHandle<double>;
8784
template class InterpolatorHandle<thrust::complex<double>>;
8885
template class InterpolatorHandle<unsigned char>;
86+
template class InterpolatorHandle<unsigned int>;
8987
} // namespace isce3::cuda::core

cxx/isce3/cuda/core/gpuBicubicInterpolator.cu

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -179,6 +179,7 @@ template class gpuBicubicInterpolator<thrust::complex<double>>;
179179
template class gpuBicubicInterpolator<float>;
180180
template class gpuBicubicInterpolator<thrust::complex<float>>;
181181
template class gpuBicubicInterpolator<unsigned char>;
182+
template class gpuBicubicInterpolator<unsigned int>;
182183

183184
template __global__ void gpuInterpolator_g<double>(
184185
gpuBicubicInterpolator<double> interp, double* x, double* y,

cxx/isce3/cuda/core/gpuBilinearInterpolator.cu

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -115,6 +115,7 @@ template class gpuBilinearInterpolator<thrust::complex<double>>;
115115
template class gpuBilinearInterpolator<float>;
116116
template class gpuBilinearInterpolator<thrust::complex<float>>;
117117
template class gpuBilinearInterpolator<unsigned char>;
118+
template class gpuBilinearInterpolator<unsigned int>;
118119

119120
template __global__ void gpuInterpolator_g<double>(
120121
gpuBilinearInterpolator<double> interp, double* x, double* y,

cxx/isce3/cuda/core/gpuNearestNeighborInterpolator.cu

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,3 +19,4 @@ template class gpuNearestNeighborInterpolator<thrust::complex<double>>;
1919
template class gpuNearestNeighborInterpolator<float>;
2020
template class gpuNearestNeighborInterpolator<thrust::complex<float>>;
2121
template class gpuNearestNeighborInterpolator<unsigned char>;
22+
template class gpuNearestNeighborInterpolator<unsigned int>;

cxx/isce3/cuda/core/gpuSpline2dInterpolator.cu

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -154,6 +154,7 @@ template class gpuSpline2dInterpolator<thrust::complex<double>>;
154154
template class gpuSpline2dInterpolator<float>;
155155
template class gpuSpline2dInterpolator<thrust::complex<float>>;
156156
template class gpuSpline2dInterpolator<unsigned char>;
157+
template class gpuSpline2dInterpolator<unsigned int>;
157158

158159
template __global__ void gpuInterpolator_g<double>(
159160
gpuSpline2dInterpolator<double> interp, double* x, double* y,

cxx/isce3/cuda/geocode/Geocode.cu

Lines changed: 29 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -201,8 +201,10 @@ Geocode::Geocode(const isce3::product::GeoGridParameters & geogrid,
201201
_interp_double_handle(data_interp_method),
202202
_interp_cdouble_handle(data_interp_method),
203203
_interp_unsigned_char_handle(data_interp_method),
204+
_interp_unsigned_int_handle(data_interp_method),
204205
_proj_handle(geogrid.epsg()),
205-
_dem_interp_method(dem_interp_method)
206+
_dem_interp_method(dem_interp_method),
207+
_data_interp_method(data_interp_method)
206208
{
207209
// init light weight radar grid
208210
_radar_grid.sensing_start = _rdr_geom.radarGrid().sensingStart();
@@ -231,10 +233,12 @@ Geocode::Geocode(const isce3::product::GeoGridParameters & geogrid,
231233
_invalid_float = std::numeric_limits<float>::quiet_NaN();
232234
_invalid_double = std::numeric_limits<double>::quiet_NaN();
233235
_invalid_unsigned_char = 255;
236+
_invalid_unsigned_int = 4294967295;
234237
} else {
235238
_invalid_float = invalid_value;
236239
_invalid_double = static_cast<double>(invalid_value);
237240
_invalid_unsigned_char = static_cast<unsigned char>(invalid_value);
241+
_invalid_unsigned_int = static_cast<unsigned int>(invalid_value);
238242
}
239243
}
240244

@@ -337,9 +341,21 @@ void Geocode::setBlockRdrCoordGrid(const size_t block_number)
337341
}
338342
}
339343

344+
void Geocode::rasterDtypeInterpCheck(const int dtype) const
345+
{
346+
if ((dtype == GDT_Byte || dtype == GDT_UInt32) &&
347+
_data_interp_method != isce3::core::NEAREST_METHOD)
348+
{
349+
std::string err_str {"int type of raster can only use nearest neighbor interp"};
350+
throw isce3::except::InvalidArgument(ISCE_SRCINFO(), err_str);
351+
}
352+
}
353+
340354
template<class T>
341355
void Geocode::geocodeRasterBlock(Raster& output_raster, Raster& input_raster)
342356
{
357+
rasterDtypeInterpCheck(input_raster.dtype());
358+
343359
// determine number of elements in output vector
344360
const auto n_elem_out = _geo_block_length * _geogrid.width();
345361

@@ -362,6 +378,9 @@ void Geocode::geocodeRasterBlock(Raster& output_raster, Raster& input_raster)
362378
} else if constexpr (std::is_same_v<T, unsigned char>) {
363379
interp = _interp_unsigned_char_handle.getInterp();
364380
invalid_value = _invalid_unsigned_char;
381+
} else if constexpr (std::is_same_v<T, unsigned int>) {
382+
interp = _interp_unsigned_int_handle.getInterp();
383+
invalid_value = _invalid_unsigned_int;
365384
}
366385

367386
// 0 width indicates current block is out of bounds
@@ -427,6 +446,10 @@ void Geocode::geocodeRasters(
427446

428447
const auto n_raster_pairs = output_rasters.size();
429448

449+
// check if raster types consistent with data interp method
450+
for (size_t i_raster = 0; i_raster < n_raster_pairs; ++i_raster)
451+
rasterDtypeInterpCheck(input_rasters[i_raster].get().dtype());
452+
430453
// iterate over blocks
431454
for (size_t i_block = 0; i_block < _n_blocks; ++i_block) {
432455

@@ -457,6 +480,10 @@ void Geocode::geocodeRasters(
457480
geocodeRasterBlock<unsigned char>(
458481
output_rasters[i_raster], input_rasters[i_raster]);
459482
break;}
483+
case GDT_UInt32: {
484+
geocodeRasterBlock<unsigned int>(
485+
output_rasters[i_raster], input_rasters[i_raster]);
486+
break;}
460487
default: {
461488
throw isce3::except::RuntimeError(ISCE_SRCINFO(),
462489
"unsupported datatype");
@@ -474,4 +501,5 @@ EXPLICIT_INSTATIATION(thrust::complex<float>);
474501
EXPLICIT_INSTATIATION(double);
475502
EXPLICIT_INSTATIATION(thrust::complex<double>);
476503
EXPLICIT_INSTATIATION(unsigned char);
504+
EXPLICIT_INSTATIATION(unsigned int);
477505
} // namespace isce3::cuda::geocode

cxx/isce3/cuda/geocode/Geocode.h

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -178,12 +178,18 @@ class Geocode {
178178
_interp_cdouble_handle;
179179
isce3::cuda::core::InterpolatorHandle<unsigned char>
180180
_interp_unsigned_char_handle;
181+
isce3::cuda::core::InterpolatorHandle<unsigned int>
182+
_interp_unsigned_int_handle;
181183

182184
// value applied to invalid geogrid pixels
183185
float _invalid_float;
184186
double _invalid_double;
185187
unsigned char _invalid_unsigned_char;
188+
unsigned int _invalid_unsigned_int;
186189

187190
isce3::core::dataInterpMethod _dem_interp_method;
191+
isce3::core::dataInterpMethod _data_interp_method;
192+
193+
void rasterDtypeInterpCheck(const int dtype) const;
188194
};
189195
} // namespace isce3::cuda::geocode

cxx/isce3/geocode/GeocodeCov.cpp

Lines changed: 33 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ void getBlocksNumberAndLength(const int array_length, const int array_width,
105105
const long long max_block_size)
106106
{
107107

108-
const int max_block_length = max_block_size /
108+
const int max_block_length = max_block_size /
109109
(nbands * array_width * type_size);
110110

111111
int _block_length = std::min(array_length, max_block_length);
@@ -290,6 +290,23 @@ void Geocode<T>::geocodeInterp(
290290
std::unique_ptr<isce3::core::ProjectionBase> proj(
291291
isce3::core::createProj(_epsgOut));
292292

293+
// make sure int type rasters only used with nearest neighbor
294+
for (int band = 0; band < nbands; ++band)
295+
{
296+
const auto dtype = inputRaster.dtype(band + 1);
297+
if ((dtype == GDT_Byte || dtype == GDT_UInt32)
298+
&& _data_interp_method != isce3::core::NEAREST_METHOD)
299+
{
300+
pyre::journal::error_t error(
301+
"isce3.geocode.GeocodeCov.geocodeInterp");
302+
error << "int type of raster can only use nearest neighbor interp"
303+
<< pyre::journal::endl;
304+
std::string err_str {
305+
"int type of raster can only use nearest neighbor interp"};
306+
throw isce3::except::InvalidArgument(ISCE_SRCINFO(), err_str);
307+
}
308+
}
309+
293310
// create data interpolator
294311
std::unique_ptr<isce3::core::Interpolator<T_out>> interp {
295312
isce3::core::createInterpolator<T_out>(_data_interp_method)};
@@ -374,7 +391,7 @@ void Geocode<T>::geocodeInterp(
374391
if (std::isnan(rtc_geogrid_upsampling))
375392
rtc_geogrid_upsampling = 1;
376393

377-
isce3::core::MemoryModeBlockY rtc_memory_mode =
394+
isce3::core::MemoryModeBlockY rtc_memory_mode =
378395
isce3::core::MemoryModeBlockY::AutoBlocksY;
379396
int radar_grid_nlooks = 1;
380397
computeRtc(demRaster, *rtc_raster, radar_grid, _orbit, _doppler,
@@ -555,9 +572,9 @@ void Geocode<T>::geocodeInterp(
555572
azimuthFirstLine = std::max(azimuthFirstLine - interp_margin, 0);
556573
rangeFirstPixel = std::max(rangeFirstPixel - interp_margin, 0);
557574

558-
azimuthLastLine = std::min(azimuthLastLine + interp_margin,
575+
azimuthLastLine = std::min(azimuthLastLine + interp_margin,
559576
static_cast<int>(radar_grid.length() - 1));
560-
rangeLastPixel = std::min(rangeLastPixel + interp_margin,
577+
rangeLastPixel = std::min(rangeLastPixel + interp_margin,
561578
static_cast<int>(radar_grid.width() - 1));
562579

563580
if (azimuthFirstLine > azimuthLastLine ||
@@ -706,7 +723,7 @@ void Geocode<T>::geocodeInterp(
706723
auto elapsed_time_milliseconds = std::chrono::duration_cast<std::chrono::milliseconds>(
707724
std::chrono::high_resolution_clock::now() - start_time);
708725
float elapsed_time = ((float) elapsed_time_milliseconds.count()) / 1e3;
709-
info << "elapsed time (GEO-IN) [s]: " << elapsed_time << pyre::journal::endl;
726+
info << "elapsed time (GEO-IN) [s]: " << elapsed_time << pyre::journal::endl;
710727
}
711728

712729
template<class T>
@@ -1048,7 +1065,7 @@ void _processUpsampledBlock(isce3::core::Matrix<T_out>* mat, size_t block,
10481065
In this case, the input type T is expected to be complex and the output
10491066
type T_out can be real or complex.
10501067
The conversion from complex (e.g. SLC) to real (SAR backscatter)
1051-
in the context of geocoding is considered as the geocoding of
1068+
in the context of geocoding is considered as the geocoding of
10521069
the covariance matrix (diagonal elements) is done by
10531070
squaring the modulus of the complex input. The conversion between
10541071
variables of same type is considered as regular geocoding and
@@ -1115,17 +1132,17 @@ void _getUpsampledBlock(
11151132
std::cout.flush();
11161133
}
11171134
auto ptr = rdrData[band]->data();
1118-
input_raster.getBlock(ptr +
1119-
block * radar_block_length * size_x,
1135+
input_raster.getBlock(ptr +
1136+
block * radar_block_length * size_x,
11201137
xidx, block * radar_block_length + yidx, size_x,
11211138
this_radar_block_length, band + 1);
11221139
}
11231140

11241141
if (radargrid_nblocks > 1) {
1125-
std::cout << "reading band " << band + 1
1142+
std::cout << "reading band " << band + 1
11261143
<< " progress: 100%" << std::endl;
11271144
}
1128-
}
1145+
}
11291146
else if (!flag_upsample_radar_grid && std::is_same<T, T_out>::value) {
11301147
/*
11311148
Enter here if:
@@ -1139,8 +1156,8 @@ void _getUpsampledBlock(
11391156
{
11401157
input_raster.getBlock(rdrData[band]->data(), xidx, yidx, size_x,
11411158
size_y, band + 1);
1142-
}
1143-
}
1159+
}
1160+
}
11441161
else if (!flag_upsample_radar_grid) {
11451162
/*
11461163
Enter here if:
@@ -1901,7 +1918,7 @@ void Geocode<T>::geocodeAreaProj(
19011918
auto elapsed_time_milliseconds = std::chrono::duration_cast<std::chrono::milliseconds>(
19021919
std::chrono::high_resolution_clock::now() - start_time);
19031920
float elapsed_time = ((float) elapsed_time_milliseconds.count()) / 1e3;
1904-
info << "elapsed time (GEO-AP) [s]: " << elapsed_time << pyre::journal::endl;
1921+
info << "elapsed time (GEO-AP) [s]: " << elapsed_time << pyre::journal::endl;
19051922
}
19061923

19071924
template<class T>
@@ -2636,7 +2653,7 @@ void Geocode<T>::_runBlock(
26362653
}
26372654

26382655
if (flag_self_intersecting_area_element) {
2639-
/*
2656+
/*
26402657
If self-intersecting, divide area element (geogrid pixel) into
26412658
two triangles and integrate them separately.
26422659
*/
@@ -2659,7 +2676,7 @@ void Geocode<T>::_runBlock(
26592676
x10_cut, size_y, size_x, w_arr_2, w_total_2, plane_orientation);
26602677
isce3::geometry::areaProjIntegrateSegment(y10_cut, y00_cut, x10_cut,
26612678
x00_cut, size_y, size_x, w_arr_2, w_total_2, plane_orientation);
2662-
2679+
26632680
w_total = 0;
26642681
/*
26652682
The new weight array `w_arr` is the sum of the absolute values of both
@@ -2882,7 +2899,7 @@ void Geocode<T>::_runBlock(
28822899
for (int band = 0; band < nbands_off_diag_terms; ++band) {
28832900
for (int i = 0; i < this_block_size_y; ++i) {
28842901
for (int j = 0; j < this_block_size_x; ++j) {
2885-
/*
2902+
/*
28862903
Since std::numeric_limits<T_out>::quiet_NaN() with
28872904
complex T_out is (or may be) undefined, we take the "real type"
28882905
if T_out (i.e. float or double) to create the NaN value and

python/extensions/pybind_isce3/cuda/geocode/cuGeocode.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,10 @@ void addbinding(pybind11::class_<Geocode>& pyGeocode)
101101
self.geocodeRasterBlock<unsigned char>(
102102
output_raster, input_raster);
103103
break;}
104+
case GDT_UInt32: {
105+
self.geocodeRasterBlock<unsigned int>(
106+
output_raster, input_raster);
107+
break;}
104108
default: {
105109
throw isce3::except::RuntimeError(ISCE_SRCINFO(),
106110
"unsupported datatype");

python/packages/nisar/workflows/h5_prep.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -746,7 +746,7 @@ def prep_ds_insar(pcfg, dst, dst_h5):
746746
units="radians", grids=grids_val,
747747
long_name='unwrapped phase')
748748
descr = f"Connected components for {pol} layer"
749-
_create_datasets(dst_h5[pol_path], igram_shape, np.uint8,
749+
_create_datasets(dst_h5[pol_path], igram_shape, np.uint32,
750750
'connectedComponents', descr=descr, units=" ",
751751
grids=grids_val,
752752
long_name='connected components')

0 commit comments

Comments
 (0)