Optimized Belief Propagation (CPU and GPU)
ProcessBpCUDA.cpp
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 #include <iostream>
28 #include "RunEval/RunEvalConstsEnums.h"
29 #include "RunImp/UtilityFuncts.h"
30 #include "BpRunProcessing/BpConstsEnumsAliases.h"
31 #include "ProcessBpCUDA.h"
32 #include "KernelBpStereo.cu"
33 
34 //return whether or not there was an error in CUDA processing
35 template<RunData_t T, unsigned int DISP_VALS, run_environment::AccSetting ACCELERATION>
36 inline run_eval::Status ProcessBpCUDA<T, DISP_VALS, ACCELERATION>::ErrorCheck(
37  const char *file, int line, bool abort) const
38 {
39  const auto code = cudaPeekAtLastError();
40  if (code != cudaSuccess) {
41  std::cout << "CUDA ERROR: " << cudaGetErrorString(code) << " " << file << " " << line << std::endl;
42  cudaGetLastError();
43  cudaDeviceReset();
44  cudaDeviceSynchronize();
45  cudaSetDevice(0);
46  if (abort) {
47  exit(code);
48  }
49  return run_eval::Status::kError;
50  }
51  return run_eval::Status::kNoError;
52 }
53 
54 //functions for processing BP to retrieve the disparity between the images
55 
56 //run the given number of iterations of BP at the current level using the given message values in global device memory
57 template<RunData_t T, unsigned int DISP_VALS, run_environment::AccSetting ACCELERATION>
58 run_eval::Status ProcessBpCUDA<T, DISP_VALS, ACCELERATION>::RunBPAtCurrentLevel(
59  const beliefprop::BpSettings& alg_settings,
60  const BpLevel<T>& current_bp_level,
61  const beliefprop::DataCostsCheckerboards<T*>& data_costs_device,
62  const beliefprop::CheckerboardMessages<T*>& messages_device,
63  T* allocated_memory) const
64 {
65  //set to prefer L1 cache since shared memory is not used in this implementation
66  cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
67  const auto kernel_thread_block_dims =
68  this->parallel_params_.OptParamsForKernel(
69  {static_cast<unsigned int>(beliefprop::BpKernel::kBpAtLevel),
70  current_bp_level.LevelProperties().level_num_});
71  const dim3 threads{kernel_thread_block_dims[0], kernel_thread_block_dims[1]};
72  //only updating half at a time
73  const dim3 grid{
74  (unsigned int)ceil((float)(current_bp_level.LevelProperties().width_checkerboard_level_) / (float)threads.x),
75  (unsigned int)ceil((float)current_bp_level.LevelProperties().height_level_ / (float)threads.y)};
76 
77  //in cuda kernel storing data one at a time (though it is coalesced), so simd_data_size not relevant here and set to 1
78  //still is a check if start of row is aligned
79  const bool data_aligned{
80  beliefprop::MemoryAlignedAtDataStart<T>(
81  0, 1, current_bp_level.LevelProperties().bytes_align_memory_,
82  current_bp_level.LevelProperties().padded_width_checkerboard_level_)};
83 
84  //at each level, run BP for numIterations, alternating between updating the messages between the two "checkerboards"
85  for (unsigned int iteration_num = 0; iteration_num < alg_settings.num_iterations; iteration_num++)
86  {
87  beliefprop::CheckerboardPart checkerboard_part_update =
88  ((iteration_num % 2) == 0) ?
89  beliefprop::CheckerboardPart::kCheckerboardPart1 :
90  beliefprop::CheckerboardPart::kCheckerboardPart0;
91  cudaDeviceSynchronize();
92 
93  using namespace beliefprop;
94  if constexpr (DISP_VALS > 0) {
95  beliefprop_cuda::RunBPIterationUsingCheckerboardUpdates<T, DISP_VALS> <<<grid, threads>>> (
96  checkerboard_part_update, current_bp_level.LevelProperties(),
97  data_costs_device[0], data_costs_device[1],
98  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
99  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
100  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
101  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
102  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
103  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
104  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
105  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
106  alg_settings.disc_k_bp, data_aligned, alg_settings.num_disp_vals);
107  }
108  else {
109  beliefprop_cuda::RunBPIterationUsingCheckerboardUpdates<T, DISP_VALS> <<<grid, threads>>> (
110  checkerboard_part_update, current_bp_level.LevelProperties(),
111  data_costs_device[0], data_costs_device[1],
112  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
113  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
114  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
115  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
116  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
117  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
118  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
119  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
120  alg_settings.disc_k_bp, data_aligned, alg_settings.num_disp_vals, allocated_memory);
121  }
122 
123  cudaDeviceSynchronize();
124  if (ErrorCheck(__FILE__, __LINE__) != run_eval::Status::kNoError) {
125  return run_eval::Status::kError;
126  }
127  }
128  return run_eval::Status::kNoError;
129 }
130 
131 //copy the computed BP message values from the current now-completed level to the corresponding slots in the next level "down" in the computation
132 //pyramid; the next level down is double the width and height of the current level so each message in the current level is copied into four "slots"
133 //in the next level down
134 //need two different "sets" of message values to avoid read-write conflicts
135 template<RunData_t T, unsigned int DISP_VALS, run_environment::AccSetting ACCELERATION>
136 run_eval::Status ProcessBpCUDA<T, DISP_VALS, ACCELERATION>::CopyMessageValuesToNextLevelDown(
137  const BpLevel<T>& current_bp_level,
138  const BpLevel<T>& next_bp_level,
139  const beliefprop::CheckerboardMessages<T*>& messages_device_copy_from,
140  const beliefprop::CheckerboardMessages<T*>& messages_device_copy_to,
141  unsigned int bp_settings_num_disp_vals) const
142 {
143  const auto kernel_thread_block_dims =
144  this->parallel_params_.OptParamsForKernel(
145  {static_cast<unsigned int>(beliefprop::BpKernel::kCopyAtLevel),
146  current_bp_level.LevelProperties().level_num_});
147  const dim3 threads{kernel_thread_block_dims[0], kernel_thread_block_dims[1]};
148  const dim3 grid{
149  (unsigned int)ceil((float)(current_bp_level.LevelProperties().width_checkerboard_level_) / (float)threads.x),
150  (unsigned int)ceil((float)(current_bp_level.LevelProperties().height_level_) / (float)threads.y)};
151 
152  cudaDeviceSynchronize();
153  if (ErrorCheck(__FILE__, __LINE__) != run_eval::Status::kNoError) {
154  return run_eval::Status::kError;
155  }
156 
157  for (const auto checkerboard_part : {beliefprop::CheckerboardPart::kCheckerboardPart0,
158  beliefprop::CheckerboardPart::kCheckerboardPart1})
159  {
160  using namespace beliefprop;
161  //call the kernel to copy the computed BP message data to the next level down in parallel in each of the two "checkerboards"
162  //storing the current message values
163  beliefprop_cuda::CopyMsgDataToNextLevel<T, DISP_VALS> <<< grid, threads >>> (
164  checkerboard_part, current_bp_level.LevelProperties(), next_bp_level.LevelProperties(),
165  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
166  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
167  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
168  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
169  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
170  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
171  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
172  messages_device_copy_from[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
173  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
174  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
175  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
176  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
177  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
178  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
179  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
180  messages_device_copy_to[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
181  bp_settings_num_disp_vals);
182 
183  cudaDeviceSynchronize();
184  if (ErrorCheck(__FILE__, __LINE__) != run_eval::Status::kNoError) {
185  return run_eval::Status::kError;
186  }
187  }
188  return run_eval::Status::kNoError;
189 }
190 
191 //initialize the data cost at each pixel with no estimated Stereo values...only the data and discontinuity costs are used
192 template<RunData_t T, unsigned int DISP_VALS, run_environment::AccSetting ACCELERATION>
193 run_eval::Status ProcessBpCUDA<T, DISP_VALS, ACCELERATION>::InitializeDataCosts(
194  const beliefprop::BpSettings& alg_settings,
195  const BpLevel<T>& current_bp_level,
196  const std::array<float*, 2>& images_target_device,
197  const beliefprop::DataCostsCheckerboards<T*>& data_costs_device) const
198 {
199  if (ErrorCheck(__FILE__, __LINE__) != run_eval::Status::kNoError) {
200  return run_eval::Status::kError;
201  }
202 
203  //since this is first kernel run in BP, set to prefer L1 cache for now since no shared memory is used by default
204  cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
205 
206  //setup execution parameters
207  //the thread size remains constant throughout but the grid size is adjusted based on the current level/kernel to run
208  const auto kernel_thread_block_dims =
209  this->parallel_params_.OptParamsForKernel(
210  {static_cast<unsigned int>(beliefprop::BpKernel::kDataCostsAtLevel),
211  0});
212  const dim3 threads{kernel_thread_block_dims[0], kernel_thread_block_dims[1]};
213  //kernel run on full-sized image to retrieve data costs at the "bottom" level of the pyramid
214  const dim3 grid{
215  (unsigned int)ceil((float)current_bp_level.LevelProperties().width_level_ / (float)threads.x),
216  (unsigned int)ceil((float)current_bp_level.LevelProperties().height_level_ / (float)threads.y)};
217 
218  //initialize the data the the "bottom" of the image pyramid
219  beliefprop_cuda::InitializeBottomLevelData<T, DISP_VALS> <<<grid, threads>>> (
220  current_bp_level.LevelProperties(),
221  images_target_device[0], images_target_device[1],
222  data_costs_device[0], data_costs_device[1],
223  alg_settings.lambda_bp, alg_settings.data_k_bp, alg_settings.num_disp_vals);
224  cudaDeviceSynchronize();
225 
226  if (ErrorCheck(__FILE__, __LINE__) != run_eval::Status::kNoError) {
227  return run_eval::Status::kError;
228  }
229 
230  return run_eval::Status::kNoError;
231 }
232 
233 //initialize the message values with no previous message values...all message values are set to 0
234 template<RunData_t T, unsigned int DISP_VALS, run_environment::AccSetting ACCELERATION>
235 run_eval::Status ProcessBpCUDA<T, DISP_VALS, ACCELERATION>::InitializeMessageValsToDefault(
236  const BpLevel<T>& current_bp_level,
237  const beliefprop::CheckerboardMessages<T*>& messages_device,
238  unsigned int bp_settings_num_disp_vals) const
239 {
240  const auto kernel_thread_block_dims =
241  this->parallel_params_.OptParamsForKernel(
242  {static_cast<unsigned int>(beliefprop::BpKernel::kInitMessageVals), 0});
243  const dim3 threads{kernel_thread_block_dims[0], kernel_thread_block_dims[1]};
244  const dim3 grid{
245  (unsigned int)ceil((float)current_bp_level.LevelProperties().width_checkerboard_level_ / (float)threads.x),
246  (unsigned int)ceil((float)current_bp_level.LevelProperties().height_level_ / (float)threads.y)};
247  using namespace beliefprop;
248 
249  //initialize all the message values for each pixel at each possible movement to the default value in the kernel
250  beliefprop_cuda::InitializeMessageValsToDefaultKernel<T, DISP_VALS> <<< grid, threads >>> (
251  current_bp_level.LevelProperties(),
252  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
253  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
254  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
255  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
256  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
257  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
258  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
259  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
260  bp_settings_num_disp_vals);
261  cudaDeviceSynchronize();
262 
263  if (ErrorCheck(__FILE__, __LINE__) != run_eval::Status::kNoError) {
264  return run_eval::Status::kError;
265  }
266 
267  return run_eval::Status::kNoError;
268 }
269 
270 template<RunData_t T, unsigned int DISP_VALS, run_environment::AccSetting ACCELERATION>
271 run_eval::Status ProcessBpCUDA<T, DISP_VALS, ACCELERATION>::InitializeDataCurrentLevel(
272  const BpLevel<T>& current_bp_level, const BpLevel<T>& prev_bp_level,
273  const beliefprop::DataCostsCheckerboards<T*>& data_costs_device,
274  const beliefprop::DataCostsCheckerboards<T*>& data_costs_device_write,
275  unsigned int bp_settings_num_disp_vals) const
276 {
277  const auto kernel_thread_block_dims =
278  this->parallel_params_.OptParamsForKernel(
279  {static_cast<unsigned int>(beliefprop::BpKernel::kDataCostsAtLevel),
280  current_bp_level.LevelProperties().level_num_});
281  const dim3 threads{kernel_thread_block_dims[0], kernel_thread_block_dims[1]};
282  //each pixel "checkerboard" is half the width of the level and there are two of them
283  //each "pixel/point" at the level belongs to one checkerboard and
284  //the four-connected neighbors of the pixel are in the counterpart checkerboard
285  const dim3 grid{
286  (unsigned int)ceil(((float)current_bp_level.LevelProperties().width_checkerboard_level_) / (float)threads.x),
287  (unsigned int)ceil((float)current_bp_level.LevelProperties().height_level_ / (float)threads.y)};
288 
289  if (ErrorCheck(__FILE__, __LINE__ ) != run_eval::Status::kNoError) {
290  return run_eval::Status::kError;
291  }
292 
293  const size_t offset_num{0};
294  for (const auto& [checkerboard_part, data_costs_write] : {
295  std::make_pair(beliefprop::CheckerboardPart::kCheckerboardPart0, data_costs_device_write[0]),
296  std::make_pair(beliefprop::CheckerboardPart::kCheckerboardPart1, data_costs_device_write[1])})
297  {
298  beliefprop_cuda::InitializeCurrentLevelData<T, DISP_VALS> <<<grid, threads>>> (
299  checkerboard_part,
300  current_bp_level.LevelProperties(), prev_bp_level.LevelProperties(),
301  data_costs_device[0], data_costs_device[1],
302  data_costs_write, ((unsigned int) offset_num / sizeof(float)),
303  bp_settings_num_disp_vals);
304  cudaDeviceSynchronize();
305  if (ErrorCheck(__FILE__, __LINE__ ) != run_eval::Status::kNoError) {
306  return run_eval::Status::kError;
307  }
308  }
309  return run_eval::Status::kNoError;
310 }
311 
312 template<RunData_t T, unsigned int DISP_VALS, run_environment::AccSetting ACCELERATION>
313 float* ProcessBpCUDA<T, DISP_VALS, ACCELERATION>::RetrieveOutputDisparity(
314  const BpLevel<T>& current_bp_level,
315  const beliefprop::DataCostsCheckerboards<T*>& data_costs_device,
316  const beliefprop::CheckerboardMessages<T*>& messages_device,
317  unsigned int bp_settings_num_disp_vals) const
318 {
319  float* result_disp_map_device;
320  cudaMalloc(
321  (void**)&result_disp_map_device,
322  current_bp_level.LevelProperties().width_level_ * current_bp_level.LevelProperties().height_level_ * sizeof(float));
323 
324  const auto kernel_thread_block_dims =
325  this->parallel_params_.OptParamsForKernel(
326  {static_cast<unsigned int>(beliefprop::BpKernel::kOutputDisp),
327  0});
328  const dim3 threads{kernel_thread_block_dims[0], kernel_thread_block_dims[1]};
329  const dim3 grid{
330  (unsigned int)ceil((float)current_bp_level.LevelProperties().width_checkerboard_level_ / (float)threads.x),
331  (unsigned int)ceil((float)current_bp_level.LevelProperties().height_level_ / (float)threads.y)};
332  using namespace beliefprop;
333 
334  beliefprop_cuda::RetrieveOutputDisparity<T, DISP_VALS> <<<grid, threads>>> (
335  current_bp_level.LevelProperties(),
336  data_costs_device[0], data_costs_device[1],
337  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
338  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
339  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
340  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart0)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
341  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesUCheckerboard)],
342  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesDCheckerboard)],
343  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesLCheckerboard)],
344  messages_device[static_cast<unsigned int>(CheckerboardPart::kCheckerboardPart1)][static_cast<unsigned int>(MessageArrays::kMessagesRCheckerboard)],
345  result_disp_map_device, bp_settings_num_disp_vals);
346  cudaDeviceSynchronize();
347  if (ErrorCheck(__FILE__, __LINE__) != run_eval::Status::kNoError) {
348  return nullptr;
349  }
350 
351  return result_disp_map_device;
352 }
353 
354 template class ProcessBpCUDA<float, 0, run_environment::AccSetting::kCUDA>;
355 template class ProcessBpCUDA<float, beliefprop::kStereoSetsToProcess[0].num_disp_vals, run_environment::AccSetting::kCUDA>;
356 template class ProcessBpCUDA<float, beliefprop::kStereoSetsToProcess[1].num_disp_vals, run_environment::AccSetting::kCUDA>;
357 template class ProcessBpCUDA<float, beliefprop::kStereoSetsToProcess[2].num_disp_vals, run_environment::AccSetting::kCUDA>;
358 template class ProcessBpCUDA<float, beliefprop::kStereoSetsToProcess[3].num_disp_vals, run_environment::AccSetting::kCUDA>;
359 template class ProcessBpCUDA<float, beliefprop::kStereoSetsToProcess[4].num_disp_vals, run_environment::AccSetting::kCUDA>;
360 template class ProcessBpCUDA<float, beliefprop::kStereoSetsToProcess[5].num_disp_vals, run_environment::AccSetting::kCUDA>;
361 template class ProcessBpCUDA<float, beliefprop::kStereoSetsToProcess[6].num_disp_vals, run_environment::AccSetting::kCUDA>;
362 template class ProcessBpCUDA<double, 0, run_environment::AccSetting::kCUDA>;
363 template class ProcessBpCUDA<double, beliefprop::kStereoSetsToProcess[0].num_disp_vals, run_environment::AccSetting::kCUDA>;
364 template class ProcessBpCUDA<double, beliefprop::kStereoSetsToProcess[1].num_disp_vals, run_environment::AccSetting::kCUDA>;
365 template class ProcessBpCUDA<double, beliefprop::kStereoSetsToProcess[2].num_disp_vals, run_environment::AccSetting::kCUDA>;
366 template class ProcessBpCUDA<double, beliefprop::kStereoSetsToProcess[3].num_disp_vals, run_environment::AccSetting::kCUDA>;
367 template class ProcessBpCUDA<double, beliefprop::kStereoSetsToProcess[4].num_disp_vals, run_environment::AccSetting::kCUDA>;
368 template class ProcessBpCUDA<double, beliefprop::kStereoSetsToProcess[5].num_disp_vals, run_environment::AccSetting::kCUDA>;
369 template class ProcessBpCUDA<double, beliefprop::kStereoSetsToProcess[6].num_disp_vals, run_environment::AccSetting::kCUDA>;
370 template class ProcessBpCUDA<halftype, 0, run_environment::AccSetting::kCUDA>;
371 template class ProcessBpCUDA<halftype, beliefprop::kStereoSetsToProcess[0].num_disp_vals, run_environment::AccSetting::kCUDA>;
372 template class ProcessBpCUDA<halftype, beliefprop::kStereoSetsToProcess[1].num_disp_vals, run_environment::AccSetting::kCUDA>;
373 template class ProcessBpCUDA<halftype, beliefprop::kStereoSetsToProcess[2].num_disp_vals, run_environment::AccSetting::kCUDA>;
374 template class ProcessBpCUDA<halftype, beliefprop::kStereoSetsToProcess[3].num_disp_vals, run_environment::AccSetting::kCUDA>;
375 template class ProcessBpCUDA<halftype, beliefprop::kStereoSetsToProcess[4].num_disp_vals, run_environment::AccSetting::kCUDA>;
376 template class ProcessBpCUDA<halftype, beliefprop::kStereoSetsToProcess[5].num_disp_vals, run_environment::AccSetting::kCUDA>;
377 template class ProcessBpCUDA<halftype, beliefprop::kStereoSetsToProcess[6].num_disp_vals, run_environment::AccSetting::kCUDA>;
BpConstsEnumsAliases.h
File with namespace for enums, constants, structures, and functions specific to belief propagation pr...
KernelBpStereo.cu
This file defines the methods to perform belief propagation for disparity map estimation from stereo ...
halftype
half halftype
Definition: RunCUDASettings.h:43
RunEvalConstsEnums.h
Contains namespace with enums and constants for implementation run evaluation.
UtilityFuncts.h
Contains namespace with utility functions for implementation.
BpLevel
Class to store and retrieve properties of a bp processing level including a data type specified as a ...
Definition: BpLevel.h:60
BpLevel::LevelProperties
const beliefprop::BpLevelProperties & LevelProperties() const
Return level properties as const reference to avoid copying and not allow it to be modified.
Definition: BpLevel.h:165
ProcessBpCUDA
Child class of ProcessBp that define functions used in processing bp in the CUDA implementation.
Definition: ProcessBpCUDA.h:48
beliefprop
Namespace for enums, constants, structures, and functions specific to belief propagation processing.
Definition: BpFileHandlingConsts.h:36
beliefprop::DataCostsCheckerboards
std::array< T, kNumCheckerboardParts > DataCostsCheckerboards
Define alias for two-element array with data costs for each bp processing checkerboard....
Definition: BpConstsEnumsAliases.h:72
beliefprop::CheckerboardPart
CheckerboardPart
Define the two checkerboard "parts" that the image is divided into.
Definition: BpConstsEnumsAliases.h:41
beliefprop::CheckerboardMessages
std::array< std::array< T, kNumMessageArrays >, kNumCheckerboardParts > CheckerboardMessages
Define alias for array with message costs for each bp processing checkerboard. Each checkerboard mes...
Definition: BpConstsEnumsAliases.h:61
beliefprop::kStereoSetsToProcess
constexpr std::array< BpStereoSet, 8 > kStereoSetsToProcess
Declare stereo sets to process with name, num disparity values, and scale factor currently conesFullS...
Definition: BpEvaluationStereoSets.h:60
run_environment::AccSetting::kCUDA
@ kCUDA
run_eval::Status
Status
Enum for status to indicate if error or no error.
Definition: RunEvalConstsEnums.h:43
run_eval::Status::kError
@ kError
run_eval::Status::kNoError
@ kNoError
beliefprop::BpLevelProperties::height_level_
unsigned int height_level_
Definition: BpLevel.h:44
beliefprop::BpLevelProperties::width_checkerboard_level_
unsigned int width_checkerboard_level_
Definition: BpLevel.h:46
beliefprop::BpLevelProperties::bytes_align_memory_
unsigned int bytes_align_memory_
Definition: BpLevel.h:45
beliefprop::BpLevelProperties::width_level_
unsigned int width_level_
Definition: BpLevel.h:43
beliefprop::BpLevelProperties::padded_width_checkerboard_level_
unsigned int padded_width_checkerboard_level_
Definition: BpLevel.h:47
beliefprop::BpSettings
Structure to store the belief propagation settings including the number of levels and iterations.
Definition: BpSettings.h:87
beliefprop::BpSettings::num_iterations
unsigned int num_iterations
Definition: BpSettings.h:90
beliefprop::BpSettings::num_disp_vals
unsigned int num_disp_vals
Number of disparity values must be set for each stereo set.
Definition: BpSettings.h:101
beliefprop::BpSettings::disc_k_bp
float disc_k_bp
Discontinuity cost cap set to high value by default but is expected to be dependent on number of disp...
Definition: BpSettings.h:98