Openholo  v4.2
Open Source Digital Holographic Library
ophRec_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 //M*/
43 #include "ophRec.h"
44 #include "ophRec_GPU.h"
45 #include "sys.h"
46 //#include "CUDA.h"
47 #include <npp.h>
48 
50 {
51  LOG("%s\n", __FUNCTION__);
52  LOG("\tMemory Allocation : ");
53  auto begin = CUR_TIME;
54  auto step = CUR_TIME;
55 
56  //CUDA *cuda = CUDA::getInstance();
57 
58  const int pnX = context_.pixel_number[_X];
59  const int pnY = context_.pixel_number[_Y];
60  const int N = pnX * pnY;
61  const int nWave = context_.waveNum;
62 
63  const Real ppX = context_.pixel_pitch[_X];
64  const Real ppY = context_.pixel_pitch[_Y];
65 
66  const Real simFrom = rec_config.SimulationFrom;
67  const Real simTo = rec_config.SimulationTo;
68  const int simStep = rec_config.SimulationStep;
69  const Real simGap = (simStep > 1) ? (simTo - simFrom) / (simStep - 1) : 0;
70 
71  const Real tx = 1 / ppX;
72  const Real ty = 1 / ppY;
73  const Real dx = tx / pnX;
74  const Real dy = ty / pnY;
75 
76  const Real htx = tx / 2;
77  const Real hty = ty / 2;
78  const Real hdx = dx / 2;
79  const Real hdy = dy / 2;
80  const Real baseX = -htx + hdx;
81  const Real baseY = -hty + hdy;
82 
83  const uint nChannel = context_.waveNum;
84 
85  RecGpuConst* device_config = nullptr;
86  HANDLE_ERROR(cudaMalloc((void**)&device_config, sizeof(RecGpuConst)));
87 
88  cuDoubleComplex *device_src;
89  cuDoubleComplex *device_dst;
90  Real *device_encode;
91  HANDLE_ERROR(cudaMalloc((void**)&device_src, N * sizeof(cuDoubleComplex)));
92  HANDLE_ERROR(cudaMalloc((void**)&device_dst, N * sizeof(cuDoubleComplex)));
93  HANDLE_ERROR(cudaMalloc((void**)&device_encode, N * sizeof(Real)));
94 
95  bool bRandomPhase = true;
96  int nThread = 512;// cuda->getMaxThreads();
97  int nBlock = (N + nThread - 1) / nThread;
98 
99  Real* max_device;
100  Real* min_device;
101  HANDLE_ERROR(cudaMalloc((void **)&max_device, sizeof(Real)));
102  HANDLE_ERROR(cudaMalloc((void**)&min_device, sizeof(Real)));
103 
104  unsigned char* nppMaxBuffer;
105  int nBuffer;
106  nppsSumGetBufferSize_64f(N, &nBuffer);
107  HANDLE_ERROR(cudaMalloc((void**)&nppMaxBuffer, nBuffer));
108 
109  LOG("%lf (s)\n", ELAPSED_TIME(step, CUR_TIME));
110 
111  for (int step = 0; step < simStep; step++)
112  {
113  Real min = MAX_DOUBLE;
114  Real max = MIN_DOUBLE;
115  for (uint ch = 0; ch < nChannel; ch++)
116  {
117  const Real lambda = context_.wave_length[ch];
118  const Real k = 2 * M_PI / lambda;
119 
120  RecGpuConst* host_config = new RecGpuConst(
121  nChannel, 1, pnX, pnY, ppX, ppY,
122  simFrom + (step * simGap), k, lambda, bRandomPhase
123  );
124 
125  HANDLE_ERROR(cudaMemcpy(device_config, host_config, sizeof(RecGpuConst), cudaMemcpyHostToDevice));
126  HANDLE_ERROR(cudaMemcpy(device_src, complex_H[ch], sizeof(Complex<Real>) * N, cudaMemcpyHostToDevice));
127 
128  cudaASMPropagation(nBlock, nThread, pnX, pnY, device_src, device_dst, device_encode, device_config);
129 
130  Real *encode = new Real[N];
131  uchar *normal = new uchar[N];
132  HANDLE_ERROR(cudaMemcpy(encode, device_encode, sizeof(Real) * N, cudaMemcpyDeviceToHost));
133 
134 
135  Real locmin, locmax;
136  nppsMax_64f(device_encode, N, max_device, nppMaxBuffer);
137  nppsMin_64f(device_encode, N, min_device, nppMaxBuffer);
138 
139  cudaMemcpy(&locmax, max_device, sizeof(Real), cudaMemcpyDeviceToHost);
140  cudaMemcpy(&locmin, min_device, sizeof(Real), cudaMemcpyDeviceToHost);
141 
142  if (min > locmin) min = locmin;
143  if (max < locmax) max = locmax;
144 
145  m_vecEncoded.push_back(encode);
146  m_vecNormalized.push_back(normal);
147 
148  delete host_config;
149  }
150 
151  LOG("step: %d => max: %e / min: %e\n", step, max, min);
152  if (nWave == 3)
153  {
154  for (int ch = 0; ch < nWave; ch++)
155  {
156  int idx = step * nWave + ch;
157  normalize(m_vecEncoded[idx], m_vecNormalized[idx], pnX, pnY, max, min);
158  }
159  }
160  else
161  normalize(m_vecEncoded[step], m_vecNormalized[step], pnX, pnY);
162  }
163 
164  HANDLE_ERROR(cudaFree(nppMaxBuffer));
165  HANDLE_ERROR(cudaFree(max_device));
166  HANDLE_ERROR(cudaFree(min_device));
167  HANDLE_ERROR(cudaFree(device_src));
168  HANDLE_ERROR(cudaFree(device_dst));
169  HANDLE_ERROR(cudaFree(device_encode));
170  HANDLE_ERROR(cudaFree(device_config));
171  LOG("Total : %lf (s)\n", ELAPSED_TIME(begin, CUR_TIME));
172 }
Real SimulationTo
Definition: ophRec.h:68
#define M_PI
Definition: define.h:52
int SimulationStep
Definition: ophRec.h:70
Real SimulationFrom
Definition: ophRec.h:69
#define MIN_DOUBLE
Definition: define.h:148
unsigned char uchar
Definition: typedef.h:64
struct KernelConst RecGpuConst
#define HANDLE_ERROR(err)
#define MAX_DOUBLE
Definition: define.h:140
float Real
Definition: typedef.h:55
def k(wvl)
Definition: Depthmap.py:16
void cudaASMPropagation(const int &nBlocks, const int &nThreads, const int &nx, const int &ny, cuDoubleComplex *src, cuDoubleComplex *dst, Real *encode, const RecGpuConst *cuda_config)
vec2 pixel_pitch
Definition: Openholo.h:101
#define _X
Definition: define.h:92
uint waveNum
Definition: Openholo.h:105
ivec2 pixel_number
Definition: Openholo.h:99
#define _Y
Definition: define.h:96
void ASM_Propagation_GPU()
Definition: ophRec_GPU.cpp:49
Complex< Real > ** complex_H
Definition: Openholo.h:489
#define ELAPSED_TIME(x, y)
Definition: function.h:59
OphConfig context_
Definition: Openholo.h:485
Real * wave_length
Definition: Openholo.h:106
unsigned int uint
Definition: typedef.h:62
#define CUR_TIME
Definition: function.h:58
void normalize(T *src, uchar *dst, int x, int y)