Openholo  v4.2
Open Source Digital Holographic Library
ophPASKernel.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 ophPASKernel.cu
48 * @brief Openholo Phase Added Stereogram with CUDA GPGPU
49 * @author Minwoo Nam
50 * @date 2023/03
51 */
52 
53 #ifndef OphPASKernel_cu__
54 #define OphPASKernel_cu__
55 
56 #include <cuda_runtime.h>
57 #include <cuda_runtime_api.h>
58 #include <cuda.h>
59 #include "typedef.h"
60 #include "ophPAS_GPU.h"
61 
62 
63 /**
64 @fn __global__ void phaseCalc(float* inRe, float* inIm, constValue val)
65 @brief CalcCompensatedPhaseÀÇ GPU¹öÀü ÇÔ¼ö
66 @return void
67 @param inRe
68 @param inIm
69 */
70 __global__
71 void cudaKernel_phaseCalc(float* inRe, float* inIm, constValue val, int c_x, int c_y, int c_z, int amplitude,
72  int sex, int sey, int sen)
73 {
74  int segy = blockIdx.x * blockDim.x + threadIdx.x;
75  int segx = blockIdx.y * blockDim.y + threadIdx.y;// coordinate in a Segment
76  int segX = sex;
77  int segY = sey;
78 
79  if ((segy < sey) && (segx < sex))
80  {
81  int segxx, segyy;
82  float theta_s, theta_c;
83  int dtheta_s, dtheta_c;
84  int idx_c, idx_s;
85  float theta;
86  int segNo = sen;
87  int tbl = 1024;
88  float amp = amplitude;
89  float pi = 3.14159265358979323846f;
90  float m2_pi = (float)(pi * 2.0);
91  float rWaveNum = 9926043.13930423;// _CGHE->rWaveNumber;
92  float R;
93  int cf_cx = val.cf_cx[segx];
94  int cf_cy = val.cf_cy[segy];
95  float xc = val.xc[segx];
96  float yc = val.yc[segy];
97  segyy = segy * segX + segx;
98  segxx = cf_cy * segNo + cf_cx;
99  R = (float)(sqrt((xc - c_x) * (xc - c_x) + (yc - c_y) * (yc - c_y) + c_z * c_z));
100  theta = rWaveNum * R;
101  theta_c = theta;
102  theta_s = theta + pi;
103  dtheta_c = ((int)(theta_c * tbl / (pi * 2.0)));
104  dtheta_s = ((int)(theta_s * tbl / (pi * 2.0)));
105  idx_c = (dtheta_c) & (tbl - 1);
106  idx_s = (dtheta_s) & (tbl - 1);
107  float costbl = val.costbl[idx_c];
108  float sintbl = val.sintbl[idx_s];
109  atomicAdd(&inRe[segyy * segNo * segNo + segxx], (float)(amplitude * costbl));
110  atomicAdd(&inIm[segyy * segNo * segNo + segxx], (float)(amplitude * sintbl));
111 
112  /*
113  inRe[segyy*sen*sen + segxx]+= (float)(amplitude * costbl);
114  inIm[segyy*sen*sen + segxx]+= (float)(amplitude * sintbl);
115  */
116  }
117 }
118 
119 extern "C"
120 {
121  void cuda_Wrapper_phaseCalc(float* inRe, float* inIm, constValue val, float& cx, float&cy,
122  float&cz, float& amp, ivec3& seg)
123  {
124  dim3 blockSize(seg[_Y] / 32 + 1, seg[_X] / 32 + 1);
125  dim3 gridSize(32, 32);
126 
127  cudaKernel_phaseCalc << <gridSize, blockSize >> > (inRe, inIm, val, cx, cy, cz, amp, seg[_X], seg[_Y], seg[_Z]);
128  }
129 }
130 
131 #endif // !OphPASKernel_cu__