Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
253 commits
Select commit Hold shift + click to select a range
ee7bb32
add FFBP first stage algorithm
May 12, 2024
2f36e1c
return pointer instead of vector to avoid copy in bindings
May 12, 2024
6cdcfe2
clang confused by auto
May 12, 2024
a361b92
add python bindings
May 12, 2024
9011c04
don't collapse to avoid extra sqrt calls
May 13, 2024
c46501d
strides are in bytes
May 13, 2024
1c32c5e
add unit test
May 13, 2024
a892394
add missing include
bhawkins May 14, 2024
23d8f84
baseband the images
bhawkins May 14, 2024
af22cf9
wip last stage of ffbp
bhawkins May 14, 2024
d5e0a15
Merge branch 'develop' into factorized_bp
bhawkins May 17, 2024
bacd5e8
initial algorithm for final stage
bhawkins May 17, 2024
ccb02bf
Python bindings for final stage
bhawkins May 17, 2024
d78de4f
add final stage to unit test
bhawkins May 17, 2024
11bd6fa
add class for 2D NFFT
bhawkins May 19, 2024
3e352e7
add NFFT2d::sizes attribute
bhawkins May 19, 2024
ffd9e9f
add some methods to help debugging
bhawkins May 19, 2024
2380081
add incomplete unit test to help debug set_spectrum
bhawkins May 19, 2024
7113d6b
fix indexing bugs in NFFT2d::set_spectrum
bhawkins May 19, 2024
f9c3cfc
Use NFFT2d to interpolate sub-images
bhawkins May 19, 2024
b67a905
reduce factor size in ffbp unit test
bhawkins May 19, 2024
3ba97aa
fix indentation and add missing openmp pragma
bhawkins May 19, 2024
2e0c902
Merge branch 'develop' into factorized_bp
bhawkins May 26, 2024
a3c1904
add factorization option to top-level workflow
bhawkins May 27, 2024
de9dd68
Fix indexing bug
bhawkins May 27, 2024
c85f379
Add debug layer for factors and tweak logging
bhawkins May 27, 2024
d0f84fb
dump progressively for the impatient
bhawkins May 27, 2024
36541cb
add option to disable periodic domain for NFFT
bhawkins May 28, 2024
ca51ab9
Don't use periodic interpolation for FBP
bhawkins May 28, 2024
9db6dcd
Merge branch 'develop' into factorized_bp
bhawkins May 30, 2024
5c3fabd
Always use nice FFT sizes for polar grid
bhawkins Jun 2, 2024
7886509
Avoid integer overflow
bhawkins Jun 2, 2024
ff7b974
wip project one subimage at a time
bhawkins Jun 2, 2024
440bf22
Make a separate function for indexing junk
bhawkins Jun 2, 2024
9085e2a
add interface to avoid allocations
bhawkins Jun 3, 2024
20b853d
make NFFTParams default constructable
bhawkins Jun 3, 2024
0f7405f
allow empty radar grids
bhawkins Jun 3, 2024
0eaee55
add python bindings
bhawkins Jun 3, 2024
4d7b98d
Merge branch 'develop' into factorized_bp
Jun 24, 2024
394e16b
Add helper for geogrid -> XYZ conversions
Jul 1, 2024
438e255
Python bindings: add polar grid ctor, fix a stride bug, and enable ar…
Jul 29, 2024
8f8bfef
Fix copy of height layer
bhawkins Aug 13, 2024
22a9c69
Decorate PolarGrid methods for CUDA compat
bhawkins Aug 13, 2024
ff8f36a
add GPU first stage FBP
bhawkins Aug 13, 2024
f3599dd
Merge remote-tracking branch 'stupo/factorized_bp' into factorized_bp
bhawkins Aug 13, 2024
2b34ecc
log timing of FBP steps
bhawkins Aug 14, 2024
f0a02da
GPU interp2d
bhawkins Aug 16, 2024
f8136b1
GPU NFFT2d
bhawkins Aug 16, 2024
07957d4
expose secret macro to cmake
bhawkins Aug 16, 2024
c81f8ea
GPU projectPolarToGeo
bhawkins Aug 16, 2024
12803d8
WIP CUDA_HOSTDEV geo2polar
bhawkins Aug 16, 2024
12d871b
Merge branch 'develop' into factorized_bp
Aug 28, 2024
ab84e1a
Merge branch 'develop' into factorized_bp
bhawkins Apr 11, 2025
ef03a01
Move polar grid setup into standalone function.
bhawkins Apr 14, 2025
c24f0e3
add PolarGrid.shape property to Python bindings
bhawkins Apr 14, 2025
6e47958
Fix ob1 error and start Doppler.
Apr 16, 2025
98530c0
Add repr method for PolarGrid
Apr 16, 2025
dc68e88
Add repr for Linspace
Apr 16, 2025
5dba196
Add polar2polar algorithm
Apr 16, 2025
a7d63b8
kludge clang openmp bug
Apr 16, 2025
4e96478
add mergePolarGrids function
Apr 16, 2025
60464f2
Merge branch 'develop' into factorized_bp
Apr 16, 2025
a4cf75b
document NFFT2d (one year later!)
Apr 16, 2025
149a717
Add helper for NFFT2d interpolation
Apr 16, 2025
08260be
Allow const input data for one-shot FFTs.
Apr 16, 2025
79336d4
Allow const input data for one-shot FFTs.
Apr 16, 2025
096c1e6
Revert "Allow const input data for one-shot FFTs."
Apr 16, 2025
f980e08
Merge branch 'const_input_fft' into factorized_bp
Apr 16, 2025
3bd3166
refactor and add NFFT2d bindings
Apr 17, 2025
1e1ab6e
avoid copy
Apr 17, 2025
5d8c148
separate NFFT from FBP final stage
Apr 17, 2025
8019d82
refactor to use new makeImageNFFT2d code
Apr 17, 2025
7d7f7c9
add look_side to PolarGrid
Apr 17, 2025
77886e8
move prototype into cpp file
Apr 17, 2025
46c3a3c
add PolarGrid.look_side to Python bindings
Apr 17, 2025
db0c1e9
save factors to HDF5
Apr 18, 2025
7b29ff0
add script for making GIF of factorization
Apr 18, 2025
c23116e
add parsers for FBP junk
Apr 18, 2025
1ef6f94
remove unnecessary ellipsoid parameter
Apr 18, 2025
d2e5b86
add algorithm for middle stages
Apr 18, 2025
49ec6d1
fix bug when padding input to makeImageNFFT2d
Apr 21, 2025
113b316
guard against excessive Doppler resolution
Apr 21, 2025
55fdab7
first version of multi-stage fbp workflow
Apr 21, 2025
502e464
eliminate redundant polar2geo calculations
Apr 22, 2025
9925e99
limit doppler resolution
Apr 22, 2025
5d96a0b
more detailed log messages
Apr 23, 2025
c2ed666
move dq calculation to separate function
Apr 23, 2025
83470dc
Use average rather than min dq and allow user to specify scaling cons…
Apr 23, 2025
a013875
fix one-left-over edge case
Apr 24, 2025
af365f2
get cuda compiling again
bhawkins Apr 24, 2025
3a3cfbb
gcc 9.4.0 openmp fixes
bhawkins Apr 24, 2025
8028226
fix formatting
bhawkins Apr 24, 2025
e7abfe9
don't copy/paste prototypes
bhawkins Apr 24, 2025
a04ba16
remove duplicate include
bhawkins Apr 24, 2025
818e3a4
repeat some refactoring for cuda
bhawkins Apr 24, 2025
db9de28
avoid overflow
bhawkins Apr 24, 2025
28c4b71
update cuda bindings
bhawkins Apr 24, 2025
d858e6e
refactor to delete redundant code
Apr 25, 2025
512170a
expose more algorithm knobs in config file
Apr 25, 2025
11aa97b
add python bindings to view NFFT2d spectrum
Apr 28, 2025
c80bab4
Add Python bindings for polar2geo_bracket
Apr 29, 2025
f927cb0
Implement bistatic corrections for polar grid
Apr 30, 2025
8f0559a
Remove densify option. Can just zero-pad
Apr 30, 2025
fd3a282
use time-domain zero-padding for fast forward FT
Apr 30, 2025
8980fd5
separate polar grid setup from backprojection
May 1, 2025
c05b222
fix output from Python bindings
May 1, 2025
456bb4e
log message improvements
May 1, 2025
a427637
include PRI in duration and time bounds
May 1, 2025
2ffd2fc
remove description of deleted densify option
May 1, 2025
de099ff
naming things
May 1, 2025
3288ce1
refactor cuda to match cpp
bhawkins May 2, 2025
a739484
fix tuple returns and make bbox return indices rather than grid
May 2, 2025
ab4eaad
update bbox python bindings to return slices
May 2, 2025
4cfaa71
fix missed header update
bhawkins May 3, 2025
966181b
refactor out some duplicated code
bhawkins May 3, 2025
ddccceb
trailing whitespace
bhawkins May 3, 2025
b23f35e
safer cast
bhawkins May 3, 2025
e49583f
fix dimension index and variable names
May 3, 2025
0633b27
check mask shape
May 3, 2025
8ce577b
avoid std::vector<bool> avoid numpy copies
May 5, 2025
43e615f
avoid using undocumented ArrayX<T> template
bhawkins May 5, 2025
c55c044
use GPU for ffbp first stage
bhawkins May 6, 2025
7097773
Split NFFT2d into two classes to reduce storage overhead
May 7, 2025
ef21f6e
Tabulating NFFT kernel doesn't make much difference in runtime, so di…
May 7, 2025
647c185
fingers getting tired
May 7, 2025
c110c27
eliminate another data buffer
May 7, 2025
0a65fc4
Add GPU NFFTKernel
bhawkins May 7, 2025
577c8c5
Refactor GPU NFFT2d to use less memory
bhawkins May 7, 2025
3aee62d
WIP cuda NFFT2d unit test
bhawkins May 8, 2025
709e1cd
Expose access to NFFT2dResult buffer to more than just friends.
bhawkins May 8, 2025
b140785
const fixes
bhawkins May 8, 2025
2ef6135
Define cast operators to copy NFFT2d result back to host.
bhawkins May 8, 2025
ebe6f41
add method to copy nfft2d result from host to device, which requires …
bhawkins May 8, 2025
7298bdb
copy/paste final stage into CUDA (was already there as a comment--oops)
bhawkins May 9, 2025
8e3aca4
Do rdr2geo on GPU
bhawkins May 12, 2025
91b47f1
do geo2rdr on GPU
bhawkins May 12, 2025
4ac3e3f
Python bindings for GPU FBP final stage.
bhawkins May 12, 2025
163de24
Fix allocation bug in NFFT2dResult CUDA copy ctor
bhawkins May 13, 2025
ab57046
Do interpolation on GPU
bhawkins May 13, 2025
e4a62e5
fix GPU phase precision :-O
bhawkins May 13, 2025
49e14ef
omg fix another bug on the same line
bhawkins May 13, 2025
2914d34
parallel stats
bhawkins May 13, 2025
85ab60a
High level Python refactor for minimal memory usage using delayed exe…
bhawkins May 15, 2025
23553f2
Make Python PolarGrid hashable (needed by previous commit)
bhawkins May 15, 2025
a3990f0
Replace thrust::copy with cudaMemcpy for complex data.
bhawkins May 15, 2025
77029ba
better syntax for pybind11 equal
bhawkins May 16, 2025
dff5c09
leave some breadcrumbs about assumptions
bhawkins May 16, 2025
ff171c7
major surgery: Use templates and iterators to avoid copying NFFT2dRes…
bhawkins May 17, 2025
caec554
Do NFFT2d on GPU
bhawkins May 18, 2025
3516715
failed to remove a copy, but clead up code a bit
bhawkins May 19, 2025
577a9b4
add subsetting for PolarGrid objects
bhawkins May 20, 2025
50a5dde
Fix off-by-one error and check assumption in gpu NFFT2d
bhawkins May 21, 2025
0fb44d8
add helper for basebanding
bhawkins May 21, 2025
4d99b7b
Implement middle stage on GPU
bhawkins May 21, 2025
18f8a65
only write FBP sub-images to file if it will persist
bhawkins May 22, 2025
53ccf86
Merge branch 'develop' into factorized_bp
bhawkins May 22, 2025
42d9d96
Merge branch 'develop' into factorized_bp
bhawkins May 22, 2025
5317c8d
fix typo in comment
bhawkins May 28, 2025
b121ce0
Implement dry tropo correction in final stage.
bhawkins Jun 3, 2025
0727717
Implement CPI mask in CUDA
bhawkins Jun 6, 2025
3c9da78
Enable OpenMP for CUDA code
bhawkins Jun 6, 2025
1c591a1
Merge branch 'develop' into factorized_bp
Jun 21, 2025
7fce6ea
Merge branch 'develop' into factorized_bp
Jun 26, 2025
7c0d973
Merge branch 'develop' into factorized_bp
bhawkins Mar 11, 2026
f85165d
change exception to warning
bhawkins Mar 11, 2026
67d6a8f
Explicitly calculate the necessary cache size.
bhawkins Mar 11, 2026
2fe3bf1
Merge branch 'develop' into factorized_bp
bhawkins Apr 6, 2026
f5a1415
Merge branch 'develop' into factorized_bp
bhawkins May 8, 2026
4b845cd
Fix possible integer overflow for large sub-images.
bhawkins May 19, 2026
1671df0
Re-initialize block data to zero before accumulating a new block.
bhawkins May 19, 2026
f907a79
Merge branch 'develop' into factorized_bp
bhawkins May 21, 2026
95f11ab
Merge branch 'develop' into factorized_bp
bhawkins May 29, 2026
9053800
Merge branch 'develop' into factorized_bp
bhawkins Jun 18, 2026
da6fc56
Merge remote-tracking branch 'origin/develop' into factorized_bp
bhawkins Jun 30, 2026
7ce059b
Use .cu extension for NFFT2d Python bindings so that thrust compiles.
bhawkins Jul 1, 2026
590eecb
Add Linspace method to return interval edges.
bhawkins Jul 23, 2026
16a9930
Add Linspace::edges unit test
bhawkins Jul 23, 2026
9c3f822
Change Linspace::edges to return std::array instead of std::pair.
bhawkins Jul 23, 2026
b15d901
Rename Linspace::edges to Linspace::bounds.
bhawkins Jul 23, 2026
2c27c56
Account for grid extents when merging them (e.g., extra half pixels).
bhawkins Jul 23, 2026
f1d725d
Fix silly bug in previous commit.
bhawkins Jul 23, 2026
1518bf7
Add bounds read-only property to Linspace bindings
bhawkins Jul 23, 2026
0a12b4f
Add Doxygen docs to PolarGrid struct members in Backproject.h
bhawkins Jul 23, 2026
1ba32d7
Add Doxygen docs to backprojectToPolarGrid function declaration
bhawkins Jul 23, 2026
cc58a02
Add Doxygen docs to accumulatePolarImagesToRadarGrid function
bhawkins Jul 23, 2026
bd19d55
Add Doxygen docs to setupPolarGridForPulses function
bhawkins Jul 23, 2026
33d471c
Add Doxygen docs to mergePolarImages function
bhawkins Jul 23, 2026
416280a
Add Doxygen docs to makeSubApertureMask function
bhawkins Jul 23, 2026
682ebd1
Improve Doxygen docs for accumulatePolarImageToGeoPoints
bhawkins Jul 23, 2026
6b1b446
Improve Doxygen docs for findPolarGridBoundingBoxInRadarCoord
bhawkins Jul 23, 2026
c315e74
Add Doxygen docs to findPolarGridBoundingBoxInRadarGrid function
bhawkins Jul 23, 2026
3f6ec5b
Add Doxygen docs to computeRadarGridGeoPoints overloads
bhawkins Jul 23, 2026
efecae0
Remove WIP comment
bhawkins Jul 23, 2026
1fadd52
Add Doxygen docs to GPU backproject functions in cuda/focus/Backproje…
bhawkins Jul 24, 2026
3015972
Add docstrings to Python bindings
bhawkins Jul 24, 2026
b5c9387
Remove trailing whitespace.
bhawkins Jul 24, 2026
0db265a
Remove unused code
bhawkins Jul 24, 2026
b475227
Add Doxygen documentation to NFFTKernel in cuda/core/Kernels.h
bhawkins Jul 30, 2026
53e31a3
Add docs to CUDA NFFT2d header
bhawkins Jul 30, 2026
c9d08b8
Rename ffbp -> fbp
bhawkins Jul 30, 2026
d65a201
Add docstrings to azcomp_bp and azcomp_fbp functions in focus.py
bhawkins Jul 30, 2026
d677ba0
Refactor azcomp_bp and azcomp_fbp into separate backprojection module
bhawkins Jul 30, 2026
fa39c22
Add docstrings to is_overlapping and Task in backprojection module
bhawkins Jul 30, 2026
45d3ae8
Fix blocking pre-PR issues in FBP branch
bhawkins Jul 30, 2026
a00b2b2
Fix non-existent error code.
bhawkins Jul 30, 2026
b356bd9
Rename backprojection.py to azcomp_bp.py
bhawkins Jul 30, 2026
ee2cf31
add module import
bhawkins Jul 30, 2026
db92e5d
Add standalone unit test for FBP
bhawkins Jul 30, 2026
58f5ce0
Add more Python bindings for NFFT2d
bhawkins Jul 31, 2026
43a8ffe
Error out for odd NFFT2d sizes
bhawkins Jul 31, 2026
491099a
Actually need to support odd sizes
bhawkins Jul 31, 2026
89f57cc
Oops, remove error for odd sizes
bhawkins Jul 31, 2026
4dc4bae
Fix typo
bhawkins Jul 31, 2026
0b264dc
Rename log unit
bhawkins Jul 31, 2026
02dddaf
Fix ob1 error that might cause unnecessary copies in GPU interp2d
bhawkins Jul 31, 2026
5dec872
Same ob1 and modulo fixes for CPU.
bhawkins Jul 31, 2026
bbab19b
Expose periodic boundary condition option to python interp1d bindings.
bhawkins Jul 31, 2026
955b74b
ob1 and modulo fix for interp1d
bhawkins Jul 31, 2026
a5648bc
Fix modulo mixing signed/unsigned operands.
bhawkins Jul 31, 2026
94d7cc5
Improve NFFT2d unit test.
bhawkins Jul 31, 2026
fa32289
Add unit test for NFFT2d python bindings
bhawkins Jul 31, 2026
77a2130
Check for degenerate n=0 case
bhawkins Jul 31, 2026
e1aca05
Revert "allow empty radar grids"
bhawkins Jul 31, 2026
5e3f235
Fix a log message
bhawkins Jul 31, 2026
733479f
Enable CUDA separable compilation for Python bindings
bhawkins Aug 3, 2026
5ca283a
Fix CUDA 13 kernel visibility for pybind linkage
bhawkins Aug 4, 2026
03d5d84
Use a concrete type for image interpolators to avoid CUDA linking iss…
bhawkins Aug 5, 2026
4482f4f
Remove deleted .icc file from install list.
bhawkins Aug 6, 2026
8342922
Remove a template parameter I missed.
bhawkins Aug 6, 2026
c77ee85
Fix dependent name template parsing ambiguity.
bhawkins Aug 6, 2026
55995b2
Revert "Fix CUDA 13 kernel visibility for pybind linkage"
bhawkins Aug 6, 2026
f71a8c9
Revert "Enable CUDA separable compilation for Python bindings"
bhawkins Aug 6, 2026
4136388
Remove dead code
bhawkins Aug 6, 2026
e99e9f5
Remove now-unnecessary helper class.
bhawkins Aug 6, 2026
3ee5684
Harmonize CPU and GPU backprojection APIs.
bhawkins Aug 6, 2026
451e9f7
Add NumPy-style docstrings to focus serialization module.
bhawkins Aug 7, 2026
37f6e94
Improve description of `writer` argument and remove redundant `rc_gri…
bhawkins Aug 7, 2026
f17eb66
Don't delay import of h5py
bhawkins Aug 7, 2026
6254d2e
Use updated azcomp API in focus.py
bhawkins Aug 7, 2026
f73b8bc
Use updated azcomp API in unit test
bhawkins Aug 7, 2026
e614079
Merge branch 'develop' into factorized_bp
bhawkins Aug 13, 2026
15754b9
Merge branch 'develop' into factorized_bp
bhawkins Aug 18, 2026
306519f
Merge branch 'develop' into factorized_bp
bhawkins Aug 25, 2026
6b6abfb
Use lower compression level for faster write speed.
bhawkins Aug 26, 2026
1cd7923
Use larger azcomp block to avoid repeated copies in 40 MHz case.
bhawkins Aug 26, 2026
08b0028
Default to new FBP algorithm
bhawkins Sep 2, 2026
86038eb
Merge remote-tracking branch 'origin/develop' into factorized_bp
bhawkins Sep 2, 2026
c5dca45
Merge remote-tracking branch 'origin/develop' into factorized_bp
bhawkins Sep 8, 2026
18a5090
Fix use-after-free risk on interpolator objects during GIL release
bhawkins Sep 10, 2026
fbb1114
Release GIL around mergePolarImages call
bhawkins Sep 10, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -90,6 +90,11 @@ if(ISCE3_WITH_CYTHON)
message(ERROR "isce3's cython extension has been removed!")
endif()

option(ISCE3_ENABLE_FBP_TIMING "Log factorized backprojection timing" OFF)
if(ISCE3_ENABLE_FBP_TIMING)
add_definitions(-DISCE3_ENABLE_FBP_TIMING)
endif()

###Layout same install directory structure as pyre
include(GNUInstallDirs)
InitInstallDirLayout()
Expand Down
1 change: 1 addition & 0 deletions cxx/isce3/Headers.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -175,6 +175,7 @@ signal/Looks.h
signal/Looks.icc
signal/multilook.h
signal/NFFT.h
signal/NFFT2d.h
signal/shiftSignal.h
signal/signalUtils.h
signal/Signal.h
Expand Down
1 change: 1 addition & 0 deletions cxx/isce3/Sources.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -103,6 +103,7 @@ signal/Filter.cpp
signal/flatten.cpp
signal/Looks.cpp
signal/NFFT.cpp
signal/NFFT2d.cpp
signal/shiftSignal.cpp
signal/signalUtils.cpp
signal/Signal.cpp
Expand Down
9 changes: 7 additions & 2 deletions cxx/isce3/core/Interp2d.icc
Original file line number Diff line number Diff line change
Expand Up @@ -33,8 +33,13 @@ DataType interp2d(const Kernel<KernelType>& kernelx,
// Do X interp at each Y index.
for (int i_kernely = 0; i_kernely < widthy; ++i_kernely) {
long i_datay = i_kernely + lowy;
if (periodic)
i_datay %= ny;
if (periodic) {
// XXX need both operands signed for correct result
const auto lny = static_cast<long>(ny);
i_datay %= lny;
// Careful that C++ modulo retains sign of dividend.
if (i_datay < 0) i_datay += lny;
}
if ((i_datay >= 0) and (i_datay < ny)) {
const DataType* zi = &z[i_datay * stridey];
const DataType* px = detail::get_contiguous_view_or_copy(
Expand Down
4 changes: 4 additions & 0 deletions cxx/isce3/core/Kernels.h
Original file line number Diff line number Diff line change
Expand Up @@ -103,6 +103,10 @@ class NFFTKernel : public Kernel<T> {

T operator()(double x) const override;

int kernel_radius() const { return _m; }
int data_size() const { return _n; }
int fft_size() const { return _fft_size; }

private:
int _m;
int _n;
Expand Down
12 changes: 12 additions & 0 deletions cxx/isce3/core/Linspace.h
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
#pragma once

#include "Common.h"
#include <array>

namespace isce3 { namespace core {

Expand Down Expand Up @@ -120,6 +121,17 @@ class Linspace {
constexpr
int search(U) const;

/**
* Get the boundaries implied by a Linspace of bin center coordinates.
*
* \returns The array [leading_edge, trailing_edge] where leading_edge is a
* half step ahead of first() and trailing_edge is a half step beyond
* last()
*/
CUDA_HOSTDEV
constexpr
std::array<T, 2> bounds() const;

private:
T _first = {};
T _spacing = {};
Expand Down
9 changes: 9 additions & 0 deletions cxx/isce3/core/Linspace.icc
Original file line number Diff line number Diff line change
Expand Up @@ -106,4 +106,13 @@ bool operator!=(const Linspace<T> & lhs, const Linspace<U> & rhs)
return !(lhs == rhs);
}

template<typename T>
CUDA_HOSTDEV
constexpr
std::array<T, 2UL> Linspace<T>::bounds() const
{
const T half = spacing() / 2;
return std::array<T, 2UL>{first() - half, last() + half};
}

}}
10 changes: 9 additions & 1 deletion cxx/isce3/core/detail/Interp1d.h
Original file line number Diff line number Diff line change
Expand Up @@ -62,8 +62,16 @@ const DataType* get_contiguous_view_or_copy(DataType block[], int width,
}
// else
if (periodic) {
// Careful that C++ modulo retains sign of dividend.
if (low < 0) {
// XXX need both operands signed for correct result
const auto lsize = static_cast<long>(size);
low %= lsize;
low += lsize;
}
for (int i = 0; i < width; ++i) {
long j = ((low + i) % size) * stride;
// Already guaranteed positive from above.
auto j = ((low + i) % size) * stride;
block[i] = data[j];
}
} else {
Expand Down
1 change: 1 addition & 0 deletions cxx/isce3/cuda/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@ target_link_libraries(${LISCECUDA} PUBLIC
${LISCE}
${CUDART_LIBRARY}
${CUDAFFT_LIBRARY}
OpenMP::OpenMP_CUDA_Optional
)

# Specify API version and build version (used to generate name and soname
Expand Down
3 changes: 3 additions & 0 deletions cxx/isce3/cuda/Headers.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,8 @@ core/gpuPoly2d.h
core/gpuProjections.h
core/Interp1d.h
core/Interp1d.icc
core/Interp2d.h
core/Interp2d.icc
core/Kernels.h
core/Kernels.icc
core/Orbit.h
Expand Down Expand Up @@ -48,4 +50,5 @@ signal/gpuCrossMul.h
signal/gpuFilter.h
signal/gpuLooks.h
signal/gpuSignal.h
signal/NFFT2d.h
)
1 change: 1 addition & 0 deletions cxx/isce3/cuda/Sources.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -53,4 +53,5 @@ signal/gpuFilter.cu
signal/gpuLooks.cu
signal/gpuRangeFilter.cu
signal/gpuSignal.cu
signal/NFFT2d.cu
)
37 changes: 37 additions & 0 deletions cxx/isce3/cuda/core/Interp2d.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
#pragma once

#include "forward.h"
#include <isce3/core/Common.h>

namespace isce3::cuda::core {

/** Interpolate Matrix z at point (x,y)
*
* @tparam KernelType kernel element type
* @tparam DataType data element type
*
* @param[in] kernelx Kernel function to use for interpolation in x direction
* @param[in] kernely Kernel function to use for interpolation in y direction
* @param[in] z Matrix to interpolate.
* @param[in] nx Number of x samples.
* @param[in] stridex Stride between x samples.
* @param[in] ny Number of y samples.
* @param[in] stridex Stride between y samples.
* @param[in] x Desired sample (0 <= x < nx).
* @param[in] y Desired sample (0 <= y < ny).
* @param[in] periodic Use periodic boundary condition. Default = false.
* @returns Interpolated value or 0 if kernel would run off array.
*
* Matrix z will be addressed as z[ix * stridex + iy * stridey] for
* 0 <= ix < nx and 0 <= iy < ny.
*/
template<typename KernelX, typename KernelY, typename DataType>
CUDA_HOSTDEV
DataType interp2d(const KernelX& kernelx,
const KernelY& kernely, const DataType* z, size_t nx,
size_t stridex, size_t ny, size_t stridey, double x, double y,
bool periodic = false);

}

#include "Interp2d.icc"
139 changes: 139 additions & 0 deletions cxx/isce3/cuda/core/Interp2d.icc
Original file line number Diff line number Diff line change
@@ -0,0 +1,139 @@
#include <isce3/math/complexOperations.h>

#include "Interp1d.h"

namespace isce3::cuda::core {

namespace detail {
template<typename Kernel, typename KT>
CUDA_HOSTDEV
void interp1d_coeffs(const Kernel& kernel, const double t,
long* low, KT coeffs[])
{
int width = int(ceil(kernel.width()));
long i0 = 0;
if (width % 2 == 0) {
i0 = static_cast<long>(ceil(t));
} else {
i0 = static_cast<long>(round(t));
}
*low = i0 - width / 2; // integer division implicit floor()
for (int i = 0; i < width; ++i) {
double ti = i + (*low) - t;
coeffs[i] = kernel(ti);
}
}

template<typename DataType>
CUDA_HOSTDEV
const DataType* get_contiguous_view_or_copy(DataType block[], int width,
long low, const DataType* data, size_t size, size_t stride,
bool periodic)
{
const long high = low + width;
if ((stride == 1) and (low >= 0) and (high <= size)) {
return &data[low];
}
// else
if (periodic) {
// Careful that C++ modulo retains sign of dividend.
if (low < 0) {
// XXX need both operands signed for correct result
const auto lsize = static_cast<long>(size);
low %= lsize;
low += lsize;
}
for (int i = 0; i < width; ++i) {
// Already guaranteed positive from above.
auto j = ((low + i) % size) * stride;
block[i] = data[j];
}
} else {
for (int i = 0; i < width; ++i) {
long j = (low + i) * stride;
if ((j >= 0) and (j < size)) {
block[i] = data[j];
} else {
block[i] = static_cast<DataType>(0);
}
}
}
return block;
}

template<typename TX, typename TY>
CUDA_HOSTDEV
auto inner_product(const int width, const TX x[], const TY y[])
{
using namespace isce3::math::complex_operations;
using TO = typename std::common_type<TX, TY>::type;
TO sum = 0;

for (int i = 0; i < width; ++i) {
sum += x[i] * y[i];
}
return sum;
}
}

template<typename KernelX, typename KernelY, typename DataType>
CUDA_HOSTDEV
DataType interp2d(const KernelX& kernelx,
const KernelY& kernely, const DataType* z, size_t nx,
size_t stridex, size_t ny, size_t stridey, double x, double y,
bool periodic)
{
using namespace isce3::math::complex_operations;

const int MAX_WIDTH = 16;

// Small-size optimization to avoid heap allocation. The fixed-size stack
// buffers below hold widthx/widthy elements, so an over-wide kernel would
// overflow them. Guard against that (device-side assert traps the kernel;
// host-side aborts) rather than silently corrupting memory.
const int widthx = static_cast<int>(ceil(kernelx.width()));
const int widthy = static_cast<int>(ceil(kernely.width()));
assert(widthx > 0 and widthx <= MAX_WIDTH);
assert(widthy > 0 and widthy <= MAX_WIDTH);

using TX = typename KernelX::value_type;
using TY = typename KernelY::value_type;
static_assert(std::is_same<TX, TY>::value);
TX coeffsx[MAX_WIDTH];
TY coeffsy[MAX_WIDTH];
DataType datax[MAX_WIDTH], datay[MAX_WIDTH];

// Pre-compute (widthx + widthy) coefficients rather than calculating them
// (widthx * widthy) times inside the loop.
// In principle we could delay calculating coeffsy until after the loop, in
// which case we'd only need a single coefficient buffer. But we need lowy
// which comes from the same API call, and we're probably okay on stack
// space anyhow.
long lowx = 0, lowy = 0;
detail::interp1d_coeffs(kernelx, x, &lowx, coeffsx);
detail::interp1d_coeffs(kernely, y, &lowy, coeffsy);

// Do X interp at each Y index.
for (int i_kernely = 0; i_kernely < widthy; ++i_kernely) {
long i_datay = i_kernely + lowy;
if (periodic) {
// XXX need both operands signed for correct result
const auto lny = static_cast<long>(ny);
i_datay %= lny;
// Careful that C++ modulo retains sign of dividend.
if (i_datay < 0) i_datay += lny;
}
if ((i_datay >= 0) and (i_datay < ny)) {
const DataType* zi = &z[i_datay * stridey];
const DataType* px = detail::get_contiguous_view_or_copy(
datax, widthx, lowx, zi, nx, stridex, periodic);
datay[i_kernely] = detail::inner_product(widthx, coeffsx, px);
} else {
datay[i_kernely] = 0;
}
}
// Do Y interp.
return detail::inner_product(widthy, coeffsy, datay);
}

}
54 changes: 54 additions & 0 deletions cxx/isce3/cuda/core/Kernels.h
Original file line number Diff line number Diff line change
Expand Up @@ -125,6 +125,60 @@ class KnabKernel : public Kernel<T, KnabKernel<T>> {
double _bandwidth;
};

/**
* NFFT time-domain kernel
*
* This is called \f$ \phi(x) \f$ in the NFFT papers @cite keiner2009,
* specifically the Kaiser-Bessel window function.
* The domain is scaled so that usage is the same as other ISCE kernels, e.g.,
* for x in [0,n) instead of [-0.5,0.5).
*/
template<typename T>
class NFFTKernel : public Kernel<T, NFFTKernel<T>> {
using Base = Kernel<T, NFFTKernel<T>>;
friend Base;

public:
/** A non-owning kernel view type that can be passed to device code */
using view_type = NFFTKernel<T>;

/**
* Construct a new NFFTKernel object.
*
* \param[in] m Half kernel size (width = 2*m+1)
* \param[in] n Length of input signal
* \param[in] fft_size FFT transform size (> n)
*/
NFFTKernel(int m, int n, int fft_size);

/** Construct from corresponding host kernel object */
NFFTKernel(const isce3::core::NFFTKernel<T>& other);

/** Get half kernel size. */
int kernel_radius() const { return m_; }

/** Get length of input signal. */
int data_size() const { return n_; }

/** Get FFT transform size. */
int fft_size() const { return fft_size_; }

explicit operator isce3::core::NFFTKernel<T>() const {
return {kernel_radius(), data_size(), fft_size()};
}

protected:
/** \internal Implementation of \p operator() */
CUDA_HOSTDEV T eval(double t) const;

private:
int m_;
int n_;
int fft_size_;
T scale_;
T b_;
};

/** A non-owning reference to a TabulatedKernel object */
template<typename T>
class TabulatedKernelView : public Kernel<T, TabulatedKernelView<T>> {
Expand Down
Loading
Loading