Openholo  v1.0
Open Source Digital Holographic Library
ophTriMesh_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 "ophTriMesh_GPU.h"
47 
48 
49 void ophTri::initialize_GPU()
50 {
51  int nx = context_.pixel_number[_X];
52  int ny = context_.pixel_number[_Y];
53  int N = nx * ny;
54 
55  if (!streamTriMesh)
56  cudaStreamCreate(&streamTriMesh);
57 
59  HANDLE_ERROR(cudaMalloc((void**)&angularSpectrum_GPU, sizeof(cufftDoubleComplex)*nx*ny));
60 
61  if (ffttemp) cudaFree(ffttemp);
62  HANDLE_ERROR(cudaMalloc((void**)&ffttemp, sizeof(cufftDoubleComplex)*nx*ny));
63 
64 }
65 void ophTri::generateAS_GPU(uint SHADING_FLAG)
66 {
68  LOG("error: WRONG SHADING_FLAG\n");
69  exit(0);
70  }
71 
72  const uint pnX = context_.pixel_number[_X];
73  const uint pnY = context_.pixel_number[_Y];
74  const uint pnXY = pnX * pnY;
75 
76  mesh_local = new Real[9];
77  Real* mesh = new Real[9];
78 
79  findNormals(SHADING_FLAG);
80 
81  HANDLE_ERROR(cudaMemsetAsync(angularSpectrum_GPU, 0, sizeof(cufftDoubleComplex) * pnXY, streamTriMesh));
82 
83  for (int j = 0; j < meshData->n_faces; j++) {
84 
85  for (int i = 0; i < 9; i++)
86  mesh[i] = scaledMeshData[9 * j + i];
87 
88  if (checkValidity(mesh, *(no + j)) != 1)
89  continue;
90 
91  if (findGeometricalRelations(mesh, *(no + j)) != 1)
92  continue;
93 
94  refAS_GPU(j, SHADING_FLAG);
95 
96  char szLog[MAX_PATH];
97  sprintf_s(szLog, "%d / %llu\n", j + 1, meshData->n_faces);
98  LOG(szLog);
99 
100  }
101 
102  HANDLE_ERROR(cudaMemsetAsync(ffttemp, 0, sizeof(cufftDoubleComplex) * pnXY, streamTriMesh));
104 
105  cufftDoubleComplex* output = (cufftDoubleComplex*)malloc(sizeof(cufftDoubleComplex) * pnXY);
106  memset(output, 0.0, sizeof(cufftDoubleComplex) * pnXY);
107 
108  HANDLE_ERROR(cudaMemcpyAsync(output, ffttemp, sizeof(cufftDoubleComplex) * pnXY, cudaMemcpyDeviceToHost), streamTriMesh);
109  //HANDLE_ERROR(cudaMemcpyAsync(output, angularSpectrum_GPU, sizeof(cufftDoubleComplex)*nx*ny, cudaMemcpyDeviceToHost), streamTriMesh);
110 
111  for (int i = 0; i < pnXY; ++i)
112  {
113  complex_H[0][i][_RE] = output[i].x;
114  complex_H[0][i][_IM] = output[i].y;
115  }
116  delete[] output;
117  delete[] mesh, scaledMeshData, no, na, nv, mesh_local;
118 }
119 
120 
121 void ophTri::refAS_GPU(int idx, uint SHADING_FLAG)
122 {
123  int nx = context_.pixel_number[_X];
124  int ny = context_.pixel_number[_Y];
125  double px = context_.pixel_pitch[_X];
126  double py = context_.pixel_pitch[_Y];
127  double waveLength = context_.wave_length[0];
128 
129  shadingFactor = 0;
130  vec3 av(0, 0, 0);
131 
132  if (SHADING_FLAG == SHADING_FLAT) {
133  vec3 no_ = no[idx];
134  n = no_ / norm(no_);
135  if (illumination[_X] == 0 && illumination[_Y] == 0 && illumination[_Z] == 0) {
136  shadingFactor = 1;
137  }
138  else {
139  vec3 normIllu = illumination / norm(illumination);
140  shadingFactor = 2 * (n[_X] * normIllu[_X] + n[_Y] * normIllu[_Y] + n[_Z] * normIllu[_Z]) + 0.3;
141  if (shadingFactor < 0)
142  shadingFactor = 0;
143  }
144  }
145  else if (SHADING_FLAG == SHADING_CONTINUOUS) {
146 
147  av[0] = nv[3 * idx + 0][0] * illumination[0] + nv[3 * idx + 0][1] * illumination[1] + nv[3 * idx + 0][2] * illumination[2] + 0.1;
148  av[2] = nv[3 * idx + 1][0] * illumination[0] + nv[3 * idx + 1][1] * illumination[1] + nv[3 * idx + 1][2] * illumination[2] + 0.1;
149  av[1] = nv[3 * idx + 2][0] * illumination[0] + nv[3 * idx + 2][1] * illumination[1] + nv[3 * idx + 2][2] * illumination[2] + 0.1;
150 
151 
152  }
153 
154  double min_double = (double)2.2250738585072014e-308;
155  double tolerence = 1e-12;
156 
157  call_cudaKernel_refAS(angularSpectrum_GPU, nx, ny, px, py, SHADING_FLAG, idx, waveLength, M_PI, shadingFactor, av[0], av[1], av[2],
158  geom.glRot[0], geom.glRot[1], geom.glRot[2], geom.glRot[3], geom.glRot[4], geom.glRot[5], geom.glRot[6], geom.glRot[7], geom.glRot[8],
159  geom.loRot[0], geom.loRot[1], geom.loRot[2], geom.loRot[3], geom.glShift[_X], geom.glShift[_Y], geom.glShift[_Z],
160  carrierWave[_X], carrierWave[_Y], carrierWave[_Z], min_double, tolerence, streamTriMesh);
161 
162 }
Real glRot[9]
Definition: ophTriMesh.h:62
Real loRot[4]
Definition: ophTriMesh.h:64
#define HANDLE_ERROR(err)
Real * wave_length
Definition: Openholo.h:69
cudaStream_t streamTriMesh
ulonglong n_faces
The number of faces in object.
Definition: ophGen.h:496
SHADING_FLAG
Definition: ophTriMesh.h:199
float Real
Definition: typedef.h:55
Real norm(const vec2 &a)
Definition: vec.h:417
void call_fftGPU(int nx, int ny, cufftDoubleComplex *input, cufftDoubleComplex *output, CUstream_st *streamTriMesh)
#define _Y
Definition: define.h:84
#define _IM
Definition: complex.h:57
#define _X
Definition: define.h:80
Real glShift[3]
Definition: ophTriMesh.h:63
void call_cudaKernel_refAS(cufftDoubleComplex *output, int nx, int ny, double px, double py, unsigned int SHADING_FLAG, int idx, double waveLength, double pi, double shadingFactor, double av0, double av1, double av2, double glRot0, double glRot1, double glRot2, double glRot3, double glRot4, double glRot5, double glRot6, double glRot7, double glRot8, double loRot0, double loRot1, double loRot2, double loRot3, double glShiftX, double glShiftY, double glShiftZ, double carrierWaveX, double carrierWaveY, double carrierWaveZ, double min_double, double tolerence, CUstream_st *streamTriMesh)
oph::ivec2 pixel_number
Definition: Openholo.h:63
#define _RE
Definition: complex.h:54
cufftDoubleComplex * angularSpectrum_GPU
OphConfig context_
Definition: Openholo.h:297
Complex< Real > ** complex_H
Definition: Openholo.h:298
#define _Z
Definition: define.h:88
unsigned int uint
Definition: typedef.h:62
#define M_PI
Definition: define.h:52
cufftDoubleComplex * ffttemp
oph::vec2 pixel_pitch
Definition: Openholo.h:64