Openholo  v1.0
Open Source Digital Holographic Library
ophDepthMap_GPU.cpp
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 #include "ophDepthMap.h"
47 #include "ophDepthMap_GPU.h"
48 #include <sys.h> //for LOG() macro
49 
50 void empty(int a, int b, int& sum)
51 {
52  int* arr;
53  int idx = 0;
54  arr = new int[b - a + 1];
55 
56  idx == b - a ? sum += arr[idx] : idx++;
57 }
58 
64 void ophDepthMap::initGPU()
65 {
66  const int nx = context_.pixel_number[0];
67  const int ny = context_.pixel_number[1];
68  const int N = nx * ny;
69 
70  if (!stream_)
71  cudaStreamCreate(&stream_);
72 
73  if (img_src_gpu) cudaFree(img_src_gpu);
74  HANDLE_ERROR(cudaMalloc((void**)&img_src_gpu, sizeof(uchar1)*N));
75 
76  if (dimg_src_gpu) cudaFree(dimg_src_gpu);
77  HANDLE_ERROR(cudaMalloc((void**)&dimg_src_gpu, sizeof(uchar1)*N));
78 
79  if (depth_index_gpu) cudaFree(depth_index_gpu);
80  if (dm_config_.FLAG_CHANGE_DEPTH_QUANTIZATION == 1)
81  HANDLE_ERROR(cudaMalloc((void**)&depth_index_gpu, sizeof(Real)*N));
82 
83  if (u_o_gpu_) cudaFree(u_o_gpu_);
84  if (u_complex_gpu_) cudaFree(u_complex_gpu_);
85 
86  HANDLE_ERROR(cudaMalloc((void**)&u_o_gpu_, sizeof(cufftDoubleComplex)*N));
87  HANDLE_ERROR(cudaMalloc((void**)&u_complex_gpu_, sizeof(cufftDoubleComplex)*N));
88 
89  if (k_temp_d_) cudaFree(k_temp_d_);
90  HANDLE_ERROR(cudaMalloc((void**)&k_temp_d_, sizeof(cufftDoubleComplex)*N));
91 }
92 
100 bool ophDepthMap::prepareInputdataGPU(uchar* imgptr, uchar* dimgptr)
101 {
102 #ifdef CHECK_PROC_TIME
103  auto begin = CUR_TIME;
104 #endif
105  uint pnX = context_.pixel_number[_X];
106  uint pnY = context_.pixel_number[_Y];
107  const ulonglong pnXY = pnX * pnY;
108 
109  if (img_src_gpu) {
110  HANDLE_ERROR(cudaMemsetAsync(img_src_gpu, 0, sizeof(uchar1)*pnXY));
111  }
112  if (dimg_src_gpu) {
113  HANDLE_ERROR(cudaMemsetAsync(dimg_src_gpu, 0, sizeof(uchar1)*pnXY));
114  }
115 
116  HANDLE_ERROR(cudaMemcpyAsync(img_src_gpu, imgptr, sizeof(uchar1)*pnXY, cudaMemcpyHostToDevice), stream_);
117  HANDLE_ERROR(cudaMemcpyAsync(dimg_src_gpu, dimgptr, sizeof(uchar1)*pnXY, cudaMemcpyHostToDevice), stream_);
118 #ifdef CHECK_PROC_TIME
119  auto end = CUR_TIME;
120  LOG("\n%s : %lf(s)\n\n", __FUNCTION__, ((std::chrono::duration<Real>)(end - begin)).count());
121 #endif
122  return true;
123 }
124 
130 void ophDepthMap::changeDepthQuanGPU()
131 {
132 #ifdef CHECK_PROC_TIME
133  auto begin = CUR_TIME;
134 #endif
135  const uint pnX = context_.pixel_number[_X];
136  const uint pnY = context_.pixel_number[_Y];
137 
138  Real temp_depth, d1, d2;
139 
140  HANDLE_ERROR(cudaMemsetAsync(depth_index_gpu, 0, sizeof(Real)*pnX*pnY, stream_));
141 
142  for (uint dtr = 0; dtr < dm_config_.num_of_depth; dtr++)
143  {
144  temp_depth = dlevel[dtr];
145  d1 = temp_depth - dstep / 2.0;
146  d2 = temp_depth + dstep / 2.0;
147 
148  cudaChangeDepthQuanKernel(stream_, pnX, pnY, depth_index_gpu, dimg_src_gpu,
149  dtr, d1, d2, dm_config_.num_of_depth, dm_config_.far_depthmap, dm_config_.near_depthmap);
150  }
151 #ifdef CHECK_PROC_TIME
152  auto end = CUR_TIME;
153  LOG("\n%s : %lf(s)\n\n", __FUNCTION__, ((std::chrono::duration<Real>)(end - begin)).count());
154 #endif
155 }
156 
170 void ophDepthMap::calcHoloGPU(void)
171 {
172 #ifdef CHECK_PROC_TIME
173  auto begin = CUR_TIME;
174 #endif
175  if (!stream_)
176  cudaStreamCreate(&stream_);
177 
178  const int pnX = context_.pixel_number[_X];
179  const int pnY = context_.pixel_number[_Y];
180  const int pnNY = pnX * pnY;
181  const int nChannel = context_.waveNum;
182 
183  size_t depth_sz = dm_config_.render_depth.size();
184 
185  for (int ch = 0; ch < nChannel; ch++) {
186  HANDLE_ERROR(cudaMemsetAsync(u_complex_gpu_, 0, sizeof(cufftDoubleComplex) * pnNY, stream_));
187  Real lambda = context_.wave_length[ch];
188  Real k = context_.k = (2 * M_PI / lambda);
189  int p;
190  for (p = 0; p < depth_sz; ++p)
191  {
192  Complex<Real> rand_phase_val;
193  getRandPhaseValue(rand_phase_val, dm_config_.RANDOM_PHASE);
194 
195  int dtr = dm_config_.render_depth[p];
196  Real temp_depth = (is_ViewingWindow) ? dlevel_transform[dtr - 1] : dlevel[dtr - 1];
197  Complex<Real> carrier_phase_delay(0, k * temp_depth);
198  carrier_phase_delay.exp();
199 
200  HANDLE_ERROR(cudaMemsetAsync(u_o_gpu_, 0, sizeof(cufftDoubleComplex) * pnNY, stream_));
201 
202  cudaDepthHoloKernel(stream_, pnX, pnY, u_o_gpu_, img_src_gpu, dimg_src_gpu, depth_index_gpu,
203  dtr, rand_phase_val[_RE], rand_phase_val[_IM], carrier_phase_delay[_RE], carrier_phase_delay[_IM], dm_config_.FLAG_CHANGE_DEPTH_QUANTIZATION, dm_config_.DEFAULT_DEPTH_QUANTIZATION);
204 
205  HANDLE_ERROR(cudaMemsetAsync(k_temp_d_, 0, sizeof(cufftDoubleComplex) * pnNY, stream_));
206  cudaFFT(stream_, pnX, pnY, u_o_gpu_, k_temp_d_, -1);
207  propagationAngularSpectrumGPU(ch, u_o_gpu_, -temp_depth);
208 
209  //LOG("Depth: %3d of %d, z = %6.5lf mm\n", dtr, dm_config_.num_of_depth, -temp_depth * 1000);
210 
211  n_percent = (int)((Real)(ch*depth_sz + p + 1) * 100 / ((Real)depth_sz * nChannel));
212  }
213 
214  cufftDoubleComplex* p_holo_gen = new cufftDoubleComplex[pnNY];
215  memset(p_holo_gen, 0, sizeof(cufftDoubleComplex) * pnNY);
216  cudaMemcpy(p_holo_gen, u_complex_gpu_, sizeof(cufftDoubleComplex) * pnNY, cudaMemcpyDeviceToHost);
217 
218  for (int n = 0; n < pnNY; n++) {
219  complex_H[ch][n][_RE] = p_holo_gen[n].x;
220  complex_H[ch][n][_IM] = p_holo_gen[n].y;
221  }
222  delete[] p_holo_gen;
223  LOG("\n%s (%d/%d) : %lf(s)\n\n", __FUNCTION__, ch + 1, nChannel, ((std::chrono::duration<Real>)(CUR_TIME - begin)).count());
224 
225  }
226 
227 #ifdef CHECK_PROC_TIME;
228  auto end = CUR_TIME;
229  LOG("\n%s : %lf(s)\n\n", __FUNCTION__, ((std::chrono::duration<Real>)(end - begin)).count());
230 #endif
231 }
232 
240 void ophDepthMap::propagationAngularSpectrumGPU(uint channel, cufftDoubleComplex* input_u, Real propagation_dist)
241 {
242  const int pnX = context_.pixel_number[_X];
243  const int pnY = context_.pixel_number[_Y];
244  const int pnXY = pnX * pnY;
245  const Real ppX = context_.pixel_pitch[_X];
246  const Real ppY = context_.pixel_pitch[_Y];
247  const Real ssX = context_.ss[_X] = pnX * ppX;
248  const Real ssY = context_.ss[_Y] = pnY * ppY;
249  Real lambda = context_.wave_length[channel];
250  Real k = context_.k = (2 * M_PI / lambda);
251 
253  ppX, ppY, ssX, ssY, lambda, context_.k, propagation_dist);
254 }
255 
257 {
258  if (u_o_gpu_) cudaFree(u_o_gpu_);
259  if (u_complex_gpu_) cudaFree(u_complex_gpu_);
260  if (k_temp_d_) cudaFree(k_temp_d_);
261 }
cufftDoubleComplex * u_o_gpu_
#define HANDLE_ERROR(err)
Real k
Definition: Openholo.h:66
uint DEFAULT_DEPTH_QUANTIZATION
default value of the depth quantization - 256
Definition: ophGen.h:481
bool RANDOM_PHASE
If true, random phase is imposed on each depth layer.
Definition: ophGen.h:485
Real * wave_length
Definition: Openholo.h:69
uint num_of_depth
Definition: ophGen.h:475
void cudaPropagation_AngularSpKernel(CUstream_st *stream_, int pnx, int pny, cufftDoubleComplex *input_d, cufftDoubleComplex *u_complex, Real ppx, Real ppy, Real ssx, Real ssy, Real lambda, Real params_k, Real propagation_dist)
Angular spectrum propagation method for GPU implementation.
unsigned char uchar
Definition: typedef.h:64
cufftDoubleComplex * u_complex_gpu_
void cudaFFT(CUstream_st *stream, int nx, int ny, cufftDoubleComplex *in_filed, cufftDoubleComplex *output_field, int direction)
Convert data from the spatial domain to the frequency domain using 2D FFT on GPU. ...
void cudaChangeDepthQuanKernel(CUstream_st *stream_, int pnx, int pny, Real *depth_index_gpu, unsigned char *dimg_src_gpu, int dtr, Real d1, Real d2, Real params_num_of_depth, Real params_far_depthmap, Real params_near_depthmap)
Quantize depth map on the GPU, only when the number of depth quantization is not the default value (i...
void free_gpu(void)
float Real
Definition: typedef.h:55
#define CUR_TIME
Definition: function.h:58
vec2 ss
Definition: Openholo.h:67
bool FLAG_CHANGE_DEPTH_QUANTIZATION
if true, change the depth quantization from the default value.
Definition: ophGen.h:479
void getRandPhaseValue(Complex< Real > &rand_phase_val, bool rand_phase)
Assign random phase value if RANDOM_PHASE == 1.
Definition: ophGen.cpp:1266
#define _Y
Definition: define.h:84
#define _IM
Definition: complex.h:57
void empty(int a, int b, int &sum)
unsigned long long ulonglong
Definition: typedef.h:67
#define _X
Definition: define.h:80
Real near_depthmap
near value of depth in object
Definition: ophGen.h:467
oph::ivec2 pixel_number
Definition: Openholo.h:63
#define _RE
Definition: complex.h:54
Real far_depthmap
far value of depth in object
Definition: ophGen.h:469
uint waveNum
Definition: Openholo.h:68
cufftDoubleComplex * k_temp_d_
void cudaDepthHoloKernel(CUstream_st *stream, int pnx, int pny, cufftDoubleComplex *u_o_gpu_, unsigned char *img_src_gpu, unsigned char *dimg_src_gpu, Real *depth_index_gpu, int dtr, Real rand_phase_val_a, Real rand_phase_val_b, Real carrier_phase_delay_a, Real carrier_phase_delay_b, int flag_change_depth_quan, unsigned int default_depth_quan)
Find each depth plane of the input image and apply carrier phase delay to it on GPU.
Openholo Point Cloud based CGH generation with CUDA GPGPU.
OphConfig context_
Definition: Openholo.h:297
Complex< Real > ** complex_H
Definition: Openholo.h:298
vector< int > render_depth
Used when only few specific depth levels are rendered, usually for test purpose.
Definition: ophGen.h:477
cudaStream_t stream_
Real sum(const vec2 &a)
Definition: vec.h:401
unsigned int uint
Definition: typedef.h:62
#define M_PI
Definition: define.h:52
oph::vec2 pixel_pitch
Definition: Openholo.h:64