Openholo  v4.2
Open Source Digital Holographic Library
ophRecKernel.cu
Go to the documentation of this file.
1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install, copy or use the software.
7 //
8 //
9 // License Agreement
10 // For Open Source Digital Holographic Library
11 //
12 // Openholo library is free software;
13 // you can redistribute it and/or modify it under the terms of the BSD 2-Clause license.
14 //
15 // Copyright (C) 2017-2024, Korea Electronics Technology Institute. All rights reserved.
16 // E-mail : contact.openholo@gmail.com
17 // Web : http://www.openholo.org
18 //
19 // Redistribution and use in source and binary forms, with or without modification,
20 // are permitted provided that the following conditions are met:
21 //
22 // 1. Redistribution's of source code must retain the above copyright notice,
23 // this list of conditions and the following disclaimer.
24 //
25 // 2. Redistribution's in binary form must reproduce the above copyright notice,
26 // this list of conditions and the following disclaimer in the documentation
27 // and/or other materials provided with the distribution.
28 //
29 // This software is provided by the copyright holders and contributors "as is" and
30 // any express or implied warranties, including, but not limited to, the implied
31 // warranties of merchantability and fitness for a particular purpose are disclaimed.
32 // In no event shall the copyright holder or contributors be liable for any direct,
33 // indirect, incidental, special, exemplary, or consequential damages
34 // (including, but not limited to, procurement of substitute goods or services;
35 // loss of use, data, or profits; or business interruption) however caused
36 // and on any theory of liability, whether in contract, strict liability,
37 // or tort (including negligence or otherwise) arising in any way out of
38 // the use of this software, even if advised of the possibility of such damage.
39 //
40 // This software contains opensource software released under GNU Generic Public License,
41 // NVDIA Software License Agreement, or CUDA supplement to Software License Agreement.
42 // Check whether software you use contains licensed software.
43 //
44 //M*/
45 
46 /**
47 * @file ophRecKernel.cu
48 * @brief Openholo Reconstruction with CUDA GPGPU
49 * @author Minwoo Nam
50 * @date 2021/12
51 */
52 
53 
54 #ifndef ophRecKernel_cu__
55 #define ophRecKernel_cu__
56 #include "ophKernel.cuh"
57 #include <cuda_runtime_api.h>
58 #include <cuda.h>
59 #include <curand_kernel.h>
60 #include <curand_uniform.h>
61 #include <device_launch_parameters.h>
62 #include "ophRec_GPU.h"
63 
64 __global__ void cudaKernel_Encode(cuDoubleComplex *src, double *dst, const RecGpuConst* config)
65 {
66  ulonglong tid = blockIdx.x * blockDim.x + threadIdx.x;
67 
68  __shared__ int pnX;
69  __shared__ int pnY;
70  __shared__ int pnXY;
71 
72  if (threadIdx.x == 0)
73  {
74  pnX = config->pnX;
75  pnY = config->pnY;
76  pnXY = pnX * pnY;
77  }
78  __syncthreads();
79 
80  if (tid < pnXY)
81  {
82  double re = src[tid].x;
83  double im = src[tid].y;
84  dst[tid] = sqrt(re * re + im * im);
85  }
86 }
87 
88 __global__ void cudaKernel_GetKernel(cuDoubleComplex *dst, const RecGpuConst* config)
89 {
90  ulonglong tid = blockIdx.x * blockDim.x + threadIdx.x;
91 
92  __shared__ int pnX;
93  __shared__ int pnY;
94  __shared__ int pnXY;
95  __shared__ double lambda;
96  __shared__ double dx;
97  __shared__ double dy;
98 
99  __shared__ double baseX;
100  __shared__ double baseY;
101 
102  __shared__ double k;
103  __shared__ double z;
104 
105  if (threadIdx.x == 0)
106  {
107  pnX = config->pnX;
108  pnY = config->pnY;
109  pnXY = pnX * pnY;
110  lambda = config->lambda;
111  dx = config->dx;
112  dy = config->dy;
113  baseX = config->baseX;
114  baseY = config->baseY;
115  k = config->k;
116  z = config->distance;
117  }
118  __syncthreads();
119 
120  if (tid < pnXY)
121  {
122  int x = tid % pnX;
123  int y = tid / pnX;
124 
125  double curX = baseX + (x * dx);
126  double curY = baseY + (y * dy);
127  double xx = curX * lambda;
128  double yy = curY * lambda;
129 
130  cuDoubleComplex kernel;
131  kernel.x = 0;
132  kernel.y = sqrt(1 - xx * xx - yy * yy) * k * z;
133  exponent_complex(&kernel);
134 
135  double re = dst[tid].x;
136  double im = dst[tid].y;
137 
138  dst[tid].x = re * kernel.x - im * kernel.y;
139  dst[tid].y = re * kernel.y + im * kernel.x;
140  }
141 }
142 
143 
144 extern "C"
145 {
146  void cudaASMPropagation(
147  const int &nBlocks, const int &nThreads, const int &nx, const int &ny,
148  cuDoubleComplex *src, cuDoubleComplex *dst, Real *encode, const RecGpuConst* cuda_config)
149  {
150  cudaFFT(nullptr, nx, ny, src, dst, CUFFT_FORWARD, false);
151 
152  cudaKernel_GetKernel << <nBlocks, nThreads >> > (dst, cuda_config);
153 
154  if (cudaDeviceSynchronize() != cudaSuccess) {
155  LOG("Cuda error: Failed to synchronize\n");
156  return;
157  }
158 
159  cudaFFT(nullptr, nx, ny, dst, src, CUFFT_INVERSE, true);
160 
161  cudaKernel_Encode << <nBlocks, nThreads >> > (src, encode, cuda_config);
162 
163 
164  }
165 }
166 
167 #endif // !OphRecKernel_cu__