Optimized Belief Propagation (CPU and GPU)
KernelBpStereoDebug.cu
Go to the documentation of this file.
1 /*
2 Copyright (C) 2024 Scott Grauer-Gray
3 
4 This program is free software; you can redistribute it and/or modify
5 it under the terms of the GNU General Public License as published by
6 the Free Software Foundation; either version 2 of the License, or
7 (at your option) any later version.
8 
9 This program is distributed in the hope that it will be useful,
10 but WITHOUT ANY WARRANTY; without even the implied warranty of
11 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
12 GNU General Public License for more details.
13 
14 You should have received a copy of the GNU General Public License
15 along with this program; if not, write to the Free Software
16 Foundation, Inc., 59 Temple Place, Suite 330, Boston, MA 02111-1307 USA
17 */
18 
27 template<RunData_t T, unsigned int DISP_VALS>
28 __global__ void beliefprop::PrintDataAndMessageValsAtPointKernel(
29  unsigned int x_val, unsigned int y_val,
30  T* data_cost_checkerboard_0, T* data_cost_checkerboard_1,
31  T* message_u_checkerboard_0, T* message_d_checkerboard_0,
32  T* message_l_checkerboard_0, T* message_r_checkerboard_0,
33  T* message_u_checkerboard_1, T* message_d_checkerboard_1,
34  T* message_l_checkerboard_1, T* message_r_checkerboard_1,
35  unsigned int width_level_checkerboard_part, unsigned int heightLevel)
36 {
37  if (((x_val + y_val) % 2) == 0) {
38  printf("x_val: %d\n", x_val);
39  printf("y_val: %d\n", y_val);
40  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
41  printf("DISP: %d\n", current_disparity);
42  printf("messageUPrevStereoCheckerboard: %f \n",
43  (float) message_u_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
44  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
45  current_disparity, DISP_VALS)]);
46  printf("messageDPrevStereoCheckerboard: %f \n",
47  (float) message_d_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
48  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
49  current_disparity, DISP_VALS)]);
50  printf("messageLPrevStereoCheckerboard: %f \n",
51  (float) message_l_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
52  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
53  current_disparity, DISP_VALS)]);
54  printf("messageRPrevStereoCheckerboard: %f \n",
55  (float) message_r_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
56  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
57  current_disparity, DISP_VALS)]);
58  printf("dataCostStereoCheckerboard: %f \n",
59  (float) data_cost_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
60  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
61  current_disparity, DISP_VALS)]);
62  }
63  } else {
64  printf("x_val: %d\n", x_val);
65  printf("y_val: %d\n", y_val);
66  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
67  printf("DISP: %d\n", current_disparity);
68  printf("messageUPrevStereoCheckerboard: %f \n",
69  (float) message_u_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
70  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
71  current_disparity, DISP_VALS)]);
72  printf("messageDPrevStereoCheckerboard: %f \n",
73  (float) message_d_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
74  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
75  current_disparity, DISP_VALS)]);
76  printf("messageLPrevStereoCheckerboard: %f \n",
77  (float) message_l_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
78  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
79  current_disparity, DISP_VALS)]);
80  printf("messageRPrevStereoCheckerboard: %f \n",
81  (float) message_r_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
82  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
83  current_disparity, DISP_VALS)]);
84  printf("dataCostStereoCheckerboard: %f \n",
85  (float) data_cost_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
86  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
87  current_disparity, DISP_VALS)]);
88  }
89  }
90 }
91 
92 
93 template<RunData_t T, unsigned int DISP_VALS>
94 __device__ void printDataAndMessageValsAtPointDevice(
95  unsigned int x_val, unsigned int y_val,
96  T* data_cost_checkerboard_0, T* data_cost_checkerboard_1,
97  T* message_u_checkerboard_0, T* message_d_checkerboard_0,
98  T* message_l_checkerboard_0, T* message_r_checkerboard_0,
99  T* message_u_checkerboard_1, T* message_d_checkerboard_1,
100  T* message_l_checkerboard_1, T* message_r_checkerboard_1,
101  unsigned int width_level_checkerboard_part, unsigned int heightLevel)
102 {
103  if (((x_val + y_val) % 2) == 0) {
104  printf("x_val: %d\n", x_val);
105  printf("y_val: %d\n", y_val);
106  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
107  printf("DISP: %d\n", current_disparity);
108  printf("messageUPrevStereoCheckerboard: %f \n",
109  (float) message_u_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
110  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
111  current_disparity, DISP_VALS)]);
112  printf("messageDPrevStereoCheckerboard: %f \n",
113  (float) message_d_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
114  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
115  current_disparity, DISP_VALS)]);
116  printf("messageLPrevStereoCheckerboard: %f \n",
117  (float) message_l_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
118  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
119  current_disparity, DISP_VALS)]);
120  printf("messageRPrevStereoCheckerboard: %f \n",
121  (float) message_r_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
122  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
123  current_disparity, DISP_VALS)]);
124  printf("dataCostStereoCheckerboard: %f \n",
125  (float) data_cost_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
126  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
127  current_disparity, DISP_VALS)]);
128  }
129  } else {
130  printf("x_val: %d\n", x_val);
131  printf("y_val: %d\n", y_val);
132  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
133  printf("DISP: %d\n", current_disparity);
134  printf("messageUPrevStereoCheckerboard: %f \n",
135  (float) message_u_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
136  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
137  current_disparity, DISP_VALS)]);
138  printf("messageDPrevStereoCheckerboard: %f \n",
139  (float) message_d_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
140  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
141  current_disparity, DISP_VALS)]);
142  printf("messageLPrevStereoCheckerboard: %f \n",
143  (float) message_l_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
144  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
145  current_disparity, DISP_VALS)]);
146  printf("messageRPrevStereoCheckerboard: %f \n",
147  (float) message_r_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
148  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
149  current_disparity, DISP_VALS)]);
150  printf("dataCostStereoCheckerboard: %f \n",
151  (float) data_cost_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
152  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
153  current_disparity, DISP_VALS)]);
154  }
155  }
156 }
157 
158 
159 template<RunData_t T, unsigned int DISP_VALS>
160 __global__ void beliefprop::PrintDataAndMessageValsToPointKernel(
161  unsigned int x_val, unsigned int y_val,
162  T* data_cost_checkerboard_0, T* data_cost_checkerboard_1,
163  T* message_u_checkerboard_0, T* message_d_checkerboard_0,
164  T* message_l_checkerboard_0, T* message_r_checkerboard_0,
165  T* message_u_checkerboard_1, T* message_d_checkerboard_1,
166  T* message_l_checkerboard_1, T* message_r_checkerboard_1,
167  unsigned int width_level_checkerboard_part, unsigned int heightLevel)
168 {
169  const unsigned int checkerboard_adjustment = (((x_val + y_val) % 2) == 0) ? ((y_val)%2) : ((y_val+1)%2);
170  if (((x_val + y_val) % 2) == 0) {
171  printf("x_val: %d\n", x_val);
172  printf("y_val: %d\n", y_val);
173  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
174  printf("DISP: %d\n", current_disparity);
175  printf("messageUPrevStereoCheckerboard: %f \n",
176  (float) message_u_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
177  x_val / 2, y_val + 1, width_level_checkerboard_part, heightLevel,
178  current_disparity, DISP_VALS)]);
179  printf("messageDPrevStereoCheckerboard: %f \n",
180  (float) message_d_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
181  x_val / 2, y_val - 1, width_level_checkerboard_part, heightLevel,
182  current_disparity, DISP_VALS)]);
183  printf("messageLPrevStereoCheckerboard: %f \n",
184  (float) message_l_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
185  x_val / 2 + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
186  current_disparity, DISP_VALS)]);
187  printf("messageRPrevStereoCheckerboard: %f \n",
188  (float) message_r_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
189  (x_val / 2 - 1) + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
190  current_disparity, DISP_VALS)]);
191  printf("dataCostStereoCheckerboard: %f \n",
192  (float) data_cost_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
193  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
194  current_disparity, DISP_VALS)]);
195  }
196  } else {
197  printf("x_val: %d\n", x_val);
198  printf("y_val: %d\n", y_val);
199  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
200  printf("DISP: %d\n", current_disparity);
201  printf("messageUPrevStereoCheckerboard: %f \n",
202  (float) message_u_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
203  x_val / 2, y_val + 1, width_level_checkerboard_part, heightLevel,
204  current_disparity, DISP_VALS)]);
205  printf("messageDPrevStereoCheckerboard: %f \n",
206  (float) message_d_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
207  x_val / 2, y_val - 1, width_level_checkerboard_part, heightLevel,
208  current_disparity, DISP_VALS)]);
209  printf("messageLPrevStereoCheckerboard: %f \n",
210  (float) message_l_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
211  x_val / 2 + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
212  current_disparity, DISP_VALS)]);
213  printf("messageRPrevStereoCheckerboard: %f \n",
214  (float) message_r_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
215  (x_val / 2 - 1) + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
216  current_disparity, DISP_VALS)]);
217  printf("dataCostStereoCheckerboard: %f \n",
218  (float) data_cost_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
219  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
220  current_disparity, DISP_VALS)]);
221  }
222  }
223 }
224 
225 
226 template<RunData_t T, unsigned int DISP_VALS>
227 __device__ void printDataAndMessageValsToPointDevice(
228  unsigned int x_val, unsigned int y_val,
229  T* data_cost_checkerboard_0, T* data_cost_checkerboard_1,
230  T* message_u_checkerboard_0, T* message_d_checkerboard_0,
231  T* message_l_checkerboard_0, T* message_r_checkerboard_0,
232  T* message_u_checkerboard_1, T* message_d_checkerboard_1,
233  T* message_l_checkerboard_1, T* message_r_checkerboard_1,
234  unsigned int width_level_checkerboard_part, unsigned int heightLevel)
235 {
236  const unsigned int checkerboard_adjustment = (((x_val + y_val) % 2) == 0) ? ((y_val)%2) : ((y_val+1)%2);
237 
238  if (((x_val + y_val) % 2) == 0) {
239  printf("x_val: %d\n", x_val);
240  printf("y_val: %d\n", y_val);
241  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
242  printf("DISP: %d\n", current_disparity);
243  printf("messageUPrevStereoCheckerboard: %f \n",
244  (float) message_u_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
245  x_val / 2, y_val + 1, width_level_checkerboard_part, heightLevel,
246  current_disparity, DISP_VALS)]);
247  printf("messageDPrevStereoCheckerboard: %f \n",
248  (float) message_d_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
249  x_val / 2, y_val - 1, width_level_checkerboard_part, heightLevel,
250  current_disparity, DISP_VALS)]);
251  printf("messageLPrevStereoCheckerboard: %f \n",
252  (float) message_l_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
253  x_val / 2 + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
254  current_disparity, DISP_VALS)]);
255  printf("messageRPrevStereoCheckerboard: %f \n",
256  (float) message_r_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
257  (x_val / 2 - 1) + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
258  current_disparity, DISP_VALS)]);
259  printf("dataCostStereoCheckerboard: %f \n",
260  (float) data_cost_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
261  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
262  current_disparity, DISP_VALS)]);
263  }
264  } else {
265  printf("x_val: %d\n", x_val);
266  printf("y_val: %d\n", y_val);
267  for (unsigned int current_disparity = 0; current_disparity < DISP_VALS; current_disparity++) {
268  printf("DISP: %d\n", current_disparity);
269  printf("messageUPrevStereoCheckerboard: %f \n",
270  (float) message_u_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
271  x_val / 2, y_val + 1, width_level_checkerboard_part, heightLevel,
272  current_disparity, DISP_VALS)]);
273  printf("messageDPrevStereoCheckerboard: %f \n",
274  (float) message_d_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
275  x_val / 2, y_val - 1, width_level_checkerboard_part, heightLevel,
276  current_disparity, DISP_VALS)]);
277  printf("messageLPrevStereoCheckerboard: %f \n",
278  (float) message_l_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
279  x_val / 2 + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
280  current_disparity, DISP_VALS)]);
281  printf("messageRPrevStereoCheckerboard: %f \n",
282  (float) message_r_checkerboard_0[beliefprop::RetrieveIndexInDataAndMessage(
283  (x_val / 2 - 1) + checkerboard_adjustment, y_val, width_level_checkerboard_part, heightLevel,
284  current_disparity, DISP_VALS)]);
285  printf("dataCostStereoCheckerboard: %f \n",
286  (float) data_cost_checkerboard_1[beliefprop::RetrieveIndexInDataAndMessage(
287  x_val / 2, y_val, width_level_checkerboard_part, heightLevel,
288  current_disparity, DISP_VALS)]);
289  }
290  }
291 }
printDataAndMessageValsToPointDevice
__device__ void printDataAndMessageValsToPointDevice(unsigned int x_val, unsigned int y_val, T *data_cost_checkerboard_0, T *data_cost_checkerboard_1, T *message_u_checkerboard_0, T *message_d_checkerboard_0, T *message_l_checkerboard_0, T *message_r_checkerboard_0, T *message_u_checkerboard_1, T *message_d_checkerboard_1, T *message_l_checkerboard_1, T *message_r_checkerboard_1, unsigned int width_level_checkerboard_part, unsigned int heightLevel)
Definition: KernelBpStereoDebug.cu:227
printDataAndMessageValsAtPointDevice
__device__ void printDataAndMessageValsAtPointDevice(unsigned int x_val, unsigned int y_val, T *data_cost_checkerboard_0, T *data_cost_checkerboard_1, T *message_u_checkerboard_0, T *message_d_checkerboard_0, T *message_l_checkerboard_0, T *message_r_checkerboard_0, T *message_u_checkerboard_1, T *message_d_checkerboard_1, T *message_l_checkerboard_1, T *message_r_checkerboard_1, unsigned int width_level_checkerboard_part, unsigned int heightLevel)
Definition: KernelBpStereoDebug.cu:94
beliefprop::PrintDataAndMessageValsAtPointKernel
ARCHITECTURE_ADDITION void PrintDataAndMessageValsAtPointKernel(unsigned int x_val, unsigned int y_val, const beliefprop::BpLevelProperties &current_bp_level, const T *data_cost_checkerboard_0, const T *data_cost_checkerboard_1, const T *message_u_checkerboard_0, const T *message_d_checkerboard_0, const T *message_l_checkerboard_0, const T *message_r_checkerboard_0, const T *message_u_checkerboard_1, const T *message_d_checkerboard_1, const T *message_l_checkerboard_1, const T *message_r_checkerboard_1, unsigned int bp_settings_disp_vals=0)
Definition: SharedBpProcessingFuncts.h:1527
beliefprop::PrintDataAndMessageValsToPointKernel
ARCHITECTURE_ADDITION void PrintDataAndMessageValsToPointKernel(unsigned int x_val, unsigned int y_val, const beliefprop::BpLevelProperties &current_bp_level, const T *data_cost_checkerboard_0, const T *data_cost_checkerboard_1, const T *message_u_checkerboard_0, const T *message_d_checkerboard_0, const T *message_l_checkerboard_0, const T *message_r_checkerboard_0, const T *message_u_checkerboard_1, const T *message_d_checkerboard_1, const T *message_l_checkerboard_1, const T *message_r_checkerboard_1, unsigned int bp_settings_disp_vals=0)
Definition: SharedBpProcessingFuncts.h:1593
beliefprop::RetrieveIndexInDataAndMessage
ARCHITECTURE_ADDITION unsigned int RetrieveIndexInDataAndMessage(unsigned int x_val, unsigned int y_val, unsigned int width, unsigned int height, unsigned int current_disparity, unsigned int total_num_disp_vals, unsigned int offset_data=0u)
Retrieve the current 1-D index value of the given point at the given disparity in the data cost and m...
Definition: SharedBpUtilFuncts.h:72