ACloudViewer  3.9.4
A Modern Library for 3D Data Processing
SparseConvTransposeBackpropFilter.h
Go to the documentation of this file.
1 // ----------------------------------------------------------------------------
2 // - CloudViewer: www.cloudViewer.org -
3 // ----------------------------------------------------------------------------
4 // Copyright (c) 2018-2024 www.cloudViewer.org
5 // SPDX-License-Identifier: MIT
6 // ----------------------------------------------------------------------------
7 
8 #pragma once
9 #include <tbb/parallel_for.h>
10 
11 #include <Eigen/Core>
12 #include <mutex>
13 
14 namespace cloudViewer {
15 namespace ml {
16 namespace impl {
17 
20 template <class TFeat,
21  class TOut,
22  class TIndex,
23  class TKernelIndex,
24  bool NORMALIZE>
26  TOut* filter_backprop,
27  const std::vector<int>& filter_dims,
28  size_t num_out,
29  const TFeat* out_importance,
30  size_t num_inp,
31  const TFeat* inp_features,
32  const TFeat* inp_neighbors_importance_sum,
33  const int64_t* inp_neighbors_row_splits,
34  const TIndex* neighbors_index,
35  const TKernelIndex* neighbors_kernel_index,
36  const TFeat* neighbors_importance,
37  const int64_t* neighbors_row_splits,
38  const TFeat* out_features_gradient) {
39  const bool NEIGHBOR_IMPORTANCE = neighbors_importance;
40 
41  const int in_channels = filter_dims[filter_dims.size() - 2];
42  const int out_channels = filter_dims[filter_dims.size() - 1];
43 
44  int num_kernel_elements = 1;
45  for (std::size_t i = 0; i < filter_dims.size() - 2; ++i) {
46  num_kernel_elements *= filter_dims[i];
47  }
48 
49  memset(filter_backprop, 0,
50  sizeof(TOut) * num_kernel_elements * in_channels * out_channels);
51  std::mutex filter_backprop_mutex;
52 
53  tbb::parallel_for(
54  tbb::blocked_range<size_t>(0, num_out, 32),
55  [&](const tbb::blocked_range<size_t>& r) {
56  int range_length = r.end() - r.begin();
57 
58  Eigen::Matrix<TFeat, Eigen::Dynamic, Eigen::Dynamic> B(
59  in_channels * num_kernel_elements, range_length);
60  B.setZero();
61  Eigen::Matrix<TFeat, Eigen::Dynamic, Eigen::Dynamic> C(
62  out_channels, range_length);
63 
64  Eigen::Array<TFeat, Eigen::Dynamic, 1> infeat(in_channels, 1);
65 
66  for (size_t out_idx = r.begin(); out_idx != r.end();
67  ++out_idx) {
68  const int out_col = out_idx - r.begin();
69  const size_t neighbor_start = neighbors_row_splits[out_idx];
70  const size_t neighbor_end =
71  neighbors_row_splits[out_idx + 1];
72 
73  C.col(out_col) = Eigen::Map<
74  const Eigen::Array<TFeat, Eigen::Dynamic, 1>>(
75  out_features_gradient + out_idx * out_channels,
76  out_channels, 1);
77 
78  for (size_t n = neighbor_start; n < neighbor_end; ++n) {
79  const size_t inp_idx = neighbors_index[n];
80  const int kernel_idx = neighbors_kernel_index[n];
81 
82  TFeat n_importance = NEIGHBOR_IMPORTANCE
83  ? neighbors_importance[n]
84  : TFeat(1);
85  for (int ic = 0; ic < in_channels; ++ic)
86  infeat(ic) =
87  inp_features[inp_idx * in_channels + ic] *
88  n_importance;
89 
90  if (NORMALIZE) {
91  TFeat normalizer(1);
92  if (NEIGHBOR_IMPORTANCE) {
93  if (inp_neighbors_importance_sum[inp_idx] !=
94  TFeat(0))
95  normalizer /= inp_neighbors_importance_sum
96  [inp_idx];
97  } else {
98  size_t num_inp_neighbors;
99  const size_t inp_neighbor_start =
100  inp_neighbors_row_splits[inp_idx];
101  const size_t inp_neighbor_end =
102  inp_neighbors_row_splits[inp_idx + 1];
103  num_inp_neighbors =
104  inp_neighbor_end - inp_neighbor_start;
105  if (num_inp_neighbors > 0)
106  normalizer /= TFeat(num_inp_neighbors);
107  }
108  for (int ic = 0; ic < in_channels; ++ic)
109  infeat(ic) *= normalizer;
110  }
111 
112  for (int ic = 0; ic < in_channels; ++ic) {
113  B(kernel_idx * in_channels + ic, out_col) +=
114  infeat(ic);
115  }
116  }
117 
118  } // out_idx
119 
120  if (out_importance) {
121  for (size_t out_idx = r.begin(); out_idx != r.end();
122  ++out_idx) {
123  const int out_col = out_idx - r.begin();
124  C.col(out_col) *= out_importance[out_idx];
125  }
126  }
127 
128  Eigen::Matrix<TOut, Eigen::Dynamic, Eigen::Dynamic> A(
129  out_channels, num_kernel_elements * in_channels);
130 
131  A = (C * B.transpose()).template cast<TOut>();
132 
133  {
134  std::lock_guard<std::mutex> lock(filter_backprop_mutex);
135  int linear_i = 0;
136  for (int j = 0; j < num_kernel_elements * in_channels; ++j)
137  for (int i = 0; i < out_channels; ++i, ++linear_i) {
138  filter_backprop[linear_i] += A(i, j);
139  }
140  }
141  });
142 }
143 
195 template <class TFeat, class TOut, class TIndex, class TKernelIndex>
197  TOut* filter_backprop,
198  const std::vector<int>& filter_dims,
199  size_t num_out,
200  const TFeat* out_importance,
201  size_t num_inp,
202  const TFeat* inp_features,
203  const TFeat* inp_neighbors_importance_sum,
204  const int64_t* inp_neighbors_row_splits,
205  const TIndex* neighbors_index,
206  const TKernelIndex* neighbors_kernel_index,
207  const TFeat* neighbors_importance,
208  const int64_t* neighbors_row_splits,
209  const TFeat* out_features_gradient,
210  bool normalize) {
211 #define FN_PARAMETERS \
212  filter_backprop, filter_dims, num_out, out_importance, num_inp, \
213  inp_features, inp_neighbors_importance_sum, \
214  inp_neighbors_row_splits, neighbors_index, neighbors_kernel_index, \
215  neighbors_importance, neighbors_row_splits, out_features_gradient
216 
217 #define CALL_TEMPLATE(NORMALIZE) \
218  if (NORMALIZE == normalize) \
219  _SparseConvTransposeBackpropFilterCPU<TFeat, TOut, TIndex, \
220  TKernelIndex, NORMALIZE>( \
221  FN_PARAMETERS);
222 
223 #define CALL_TEMPLATE2 \
224  CALL_TEMPLATE(true) \
225  CALL_TEMPLATE(false)
226 
228 
229 #undef CALL_TEMPLATE
230 #undef CALL_TEMPLATE2
231 
232 #undef FN_PARAMETERS
233 }
234 
235 } // namespace impl
236 } // namespace ml
237 } // namespace cloudViewer
#define CALL_TEMPLATE2
__host__ __device__ float2 normalize(float2 v)
Definition: cutil_math.h:1179
void SparseConvTransposeBackpropFilterCPU(TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TFeat *out_importance, size_t num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient, bool normalize)
void _SparseConvTransposeBackpropFilterCPU(TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TFeat *out_importance, size_t num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient)
Generic file read and write utility for python interface.