LSSTApplications  10.0-2-g4f67435,11.0.rc2+1,11.0.rc2+12,11.0.rc2+3,11.0.rc2+4,11.0.rc2+5,11.0.rc2+6,11.0.rc2+7,11.0.rc2+8
LSSTDataManagementBasePackage
CudaLanczos.h
Go to the documentation of this file.
1 // -*- LSST-C++ -*-
2 
3 /*
4  * LSST Data Management System
5  * Copyright 2008 - 2012 LSST Corporation.
6  *
7  * This product includes software developed by the
8  * LSST Project (http://www.lsst.org/).
9  *
10  * This program is free software: you can redistribute it and/or modify
11  * it under the terms of the GNU General Public License as published by
12  * the Free Software Foundation, either version 3 of the License, or
13  * (at your option) any later version.
14  *
15  * This program is distributed in the hope that it will be useful,
16  * but WITHOUT ANY WARRANTY; without even the implied warranty of
17  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
18  * GNU General Public License for more details.
19  *
20  * You should have received a copy of the LSST License Statement and
21  * the GNU General Public License along with this program. If not,
22  * see <http://www.lsstcorp.org/LegalNotices/>.
23  */
24 
36 #ifdef NVCC_COMPILING
37  #define CPU_GPU __device__ __host__
38 #else
39  #define CPU_GPU
40 #endif
41 
42 namespace lsst {
43 namespace afw {
44 namespace math {
45 namespace detail {
46 namespace gpu {
47 
50 
51 int const SIZE_X_WARPING_BLOCK=16;
52 int const SIZE_Y_WARPING_BLOCK=16;
54 
56 struct SPoint2
57 {
58  double x;
59  double y;
60 
61  CPU_GPU SPoint2(double par_x, double par_y) : x(par_x), y(par_y) {}
62 
63  #ifndef NVCC_COMPILING
64  SPoint2(lsst::afw::geom::Point2D p) : x(p.getX()), y(p.getY()) {}
65  #endif
66 };
67 
69 struct SVec2
70 {
71  double x;
72  double y;
73 
74  CPU_GPU SVec2(double par_x, double par_y) : x(par_x), y(par_y) {}
75  CPU_GPU SVec2(SPoint2 a, SPoint2 b) : x(b.x-a.x), y(b.y-a.y) {}
76 
77  #ifndef NVCC_COMPILING
78  SVec2(lsst::afw::geom::Extent2D e) : x(e.getX()), y(e.getY()) {}
79  #endif
80 };
81 
83 {
84  return SVec2(a.x+b.x, a.y+b.y);
85 }
87 {
88  return SVec2(a.x-b.x, a.y-b.y);
89 }
90 CPU_GPU inline SVec2 VecMul(SVec2 v, double m)
91 {
92  return SVec2(m*v.x, m*v.y);
93 }
95 {
96  return SPoint2(p.x+v.x, p.y+v.y);
97 }
98 
99 
101 struct SBox2I
102 {
103  int begX;
104  int begY;
105  int endX;
106  int endY;
107 
108  SBox2I() {};
109 
110  CPU_GPU SBox2I(int par_begX, int par_begY, int par_endX, int par_endY)
111  : begX(par_begX), begY(par_begY), endX(par_endX), endY(par_endY) {}
112 
114  {
115  return begX <= p.x && p.x < endX
116  && begY <= p.y && p.y < endY;
117  }
118 };
119 
128 {
131 
132  CPU_GPU LinearInterp(SPoint2 par_o, SVec2 par_deltaX) : o(par_o), deltaX(par_deltaX) {};
133 
136  {
137  return MovePoint(o, VecMul(deltaX,subX) );
138  }
139 };
140 
141 
150 {
153  SVec2 ddX;
156 
157  BilinearInterp() : o(0,0), d0X(0,0), ddX(0,0), deltaY(0,0) {};
158 
159  CPU_GPU BilinearInterp(SPoint2 par_o, SVec2 par_d0X, SVec2 par_ddX, SVec2 par_deltaY)
160  : o(par_o), d0X(par_d0X), ddX(par_ddX), deltaY(par_deltaY) {}
161 
162 
165  {
166  SVec2 deltaX=VecAdd(d0X, VecMul(ddX, subY));
167  SPoint2 lineBeg= MovePoint(o, VecMul(deltaY,subY) );
168  return LinearInterp(lineBeg, deltaX);
169  }
170 
172  CPU_GPU SPoint2 Interpolate(int subX, int subY)
173  {
174  LinearInterp lineY=GetLinearInterp(subY);
175  return lineY.Interpolate(subX);
176  }
177 };
178 
180 template<typename T>
181 struct PixelIVM
182 {
183  T img;
186 };
187 
190 
192 template<typename T>
194 {
195  T* img;
201  int width;
202  int height;
203 };
204 
219 template<typename DestPixelT, typename SrcPixelT>
220 void WarpImageGpuCallKernel(bool isMaskedImage,
221  ImageDataPtr<DestPixelT> destImageGpu,
222  ImageDataPtr<SrcPixelT> srcImageGpu,
223  int mainKernelSize,
224  KernelType maskKernelType,
225  int maskKernelSize,
226  SBox2I srcGoodBox,
228  BilinearInterp* srcPosInterp,
229  int interpLength
230  );
231 
232 }}}}} //namespace lsst::afw::math::detail::gpu ends
SVec2 d0X
defines the value at origin
Definition: CudaLanczos.h:152
boost::uint16_t MaskPixel
SPoint2 MovePoint(SPoint2 p, SVec2 v)
Definition: CudaLanczos.h:94
SPoint2(lsst::afw::geom::Point2D p)
Definition: CudaLanczos.h:64
SVec2(SPoint2 a, SPoint2 b)
Definition: CudaLanczos.h:75
lsst::afw::image::MaskPixel MskPixel
Definition: CudaLanczos.h:49
SVec2 deltaX
defines the value at the origin
Definition: CudaLanczos.h:130
SVec2 VecMul(SVec2 v, double m)
Definition: CudaLanczos.h:90
SVec2(lsst::afw::geom::Extent2D e)
Definition: CudaLanczos.h:78
defines memory region containing image data
Definition: CudaLanczos.h:193
void WarpImageGpuCallKernel(bool isMaskedImage, ImageDataPtr< DestPixelT > destImageGpu, ImageDataPtr< SrcPixelT > srcImageGpu, int mainKernelSize, KernelType maskKernelType, int maskKernelSize, SBox2I srcGoodBox, PixelIVM< DestPixelT > edgePixel, BilinearInterp *srcPosInterp, int interpLength)
Calls the GPU kernel for lanczos resampling.
int const SIZE_X_WARPING_BLOCK
Definition: CudaLanczos.h:51
defines a 2D range of integer values begX &lt;= x &lt; endX, begY &lt;= y &lt; endY
Definition: CudaLanczos.h:101
SPoint2(double par_x, double par_y)
Definition: CudaLanczos.h:61
int const SIZE_MAX_WARPING_KERNEL
Definition: CudaLanczos.h:53
bool isInsideBox(gpu::SPoint2 p)
Definition: CudaLanczos.h:113
LinearInterp GetLinearInterp(int subY)
intersects the interpolation surface with a const-y plane
Definition: CudaLanczos.h:164
Simple 2D point (suitable for use on a GPU)
Definition: CudaLanczos.h:56
SPoint2 Interpolate(int subX)
Calculates a value of the interpolation function at subX.
Definition: CudaLanczos.h:135
SVec2 VecAdd(SVec2 a, SVec2 b)
Definition: CudaLanczos.h:82
#define CPU_GPU
Definition: CudaLanczos.h:39
lsst::afw::image::VariancePixel VarPixel
Definition: CudaLanczos.h:48
ImageT::SinglePixel edgePixel(lsst::afw::image::detail::Image_tag)
Return an off-the-edge pixel appropriate for a given Image type.
defines a pixel having image, variance and mask planes
Definition: CudaLanczos.h:181
float VariancePixel
! default type for Masks and MaskedImage Masks
SPoint2 Interpolate(int subX, int subY)
Calculates a value of the interpolation surface at a point (subX,subY)
Definition: CudaLanczos.h:172
BilinearInterp()
difference of neighboring values in the first column (gradient of a line at x=0)
Definition: CudaLanczos.h:157
Simple 2D vector (suitable for use on a GPU)
Definition: CudaLanczos.h:69
tuple m
Definition: lsstimport.py:48
SBox2I(int par_begX, int par_begY, int par_endX, int par_endY)
Definition: CudaLanczos.h:110
int const SIZE_Y_WARPING_BLOCK
Definition: CudaLanczos.h:52
afw::table::Key< double > b
SVec2 VecSub(SVec2 a, SVec2 b)
Definition: CudaLanczos.h:86
SVec2(double par_x, double par_y)
Definition: CudaLanczos.h:74
BilinearInterp(SPoint2 par_o, SVec2 par_d0X, SVec2 par_ddX, SVec2 par_deltaY)
Definition: CudaLanczos.h:159
LinearInterp(SPoint2 par_o, SVec2 par_deltaX)
difference of neighboring values of the function (the gradient)
Definition: CudaLanczos.h:132
SVec2 ddX
difference of difference of neighboring values in two neighbouring rows (diff. of gradients at x=0 an...
Definition: CudaLanczos.h:154