ACloudViewer  3.9.4
A Modern Library for 3D Data Processing
SparseConvTranspose.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 
13 namespace cloudViewer {
14 namespace ml {
15 namespace impl {
16 
19 template <class TFeat,
20  class TOut,
21  class TIndex,
22  class TKernelIndex,
23  bool NORMALIZE>
25  TOut* out_features,
26  const std::vector<int>& filter_dims,
27  const TFeat* filter,
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* neighbor_index,
35  const TKernelIndex* neighbors_kernel_index,
36  const TFeat* neighbor_importance,
37  const int64_t* neighbors_row_splits) {
38  const bool NEIGHBOR_IMPORTANCE = inp_neighbors_importance_sum;
39 
40  const int in_channels = filter_dims[filter_dims.size() - 2];
41  const int out_channels = filter_dims[filter_dims.size() - 1];
42 
43  int num_kernel_elements = 1;
44  for (std::size_t i = 0; i < filter_dims.size() - 2; ++i) {
45  num_kernel_elements *= filter_dims[i];
46  }
47 
48  memset(out_features, 0, sizeof(TOut) * num_out * out_channels);
49 
50  tbb::parallel_for(
51  tbb::blocked_range<size_t>(0, num_out, 32),
52  [&](const tbb::blocked_range<size_t>& r) {
53  int range_length = r.end() - r.begin();
54 
55  Eigen::Map<Eigen::Matrix<TOut, Eigen::Dynamic, Eigen::Dynamic>>
56  C(out_features + (r.begin() * out_channels),
57  out_channels, range_length);
58 
59  for (size_t out_idx = r.begin(); out_idx != r.end();
60  ++out_idx) {
61  const int out_col = out_idx - r.begin();
62  const size_t neighbor_start = neighbors_row_splits[out_idx];
63  const size_t neighbor_end =
64  neighbors_row_splits[out_idx + 1];
65 
66  for (size_t n = neighbor_start; n < neighbor_end; ++n) {
67  const size_t inp_idx = neighbor_index[n];
68  const int kernel_idx = neighbors_kernel_index[n];
69 
70  TFeat n_importance = NEIGHBOR_IMPORTANCE
71  ? neighbor_importance[n]
72  : TFeat(1);
73 
74  TFeat normalizer(1);
75  if (NORMALIZE) {
76  if (NEIGHBOR_IMPORTANCE) {
77  if (inp_neighbors_importance_sum[inp_idx] !=
78  TFeat(0))
79  normalizer /= inp_neighbors_importance_sum
80  [inp_idx];
81  } else {
82  size_t num_inp_neighbors;
83  const size_t inp_neighbor_start =
84  inp_neighbors_row_splits[inp_idx];
85  const size_t inp_neighbor_end =
86  inp_neighbors_row_splits[inp_idx + 1];
87  num_inp_neighbors =
88  inp_neighbor_end - inp_neighbor_start;
89  if (num_inp_neighbors > 0)
90  normalizer /= TFeat(num_inp_neighbors);
91  }
92  }
93 
94  Eigen::Map<const Eigen::Matrix<TFeat, Eigen::Dynamic,
95  Eigen::Dynamic>>
96  A(filter + kernel_idx * out_channels *
97  in_channels,
98  out_channels, in_channels);
99 
100  Eigen::Map<const Eigen::Matrix<TFeat, Eigen::Dynamic,
101  Eigen::Dynamic>>
102  B(inp_features + inp_idx * in_channels,
103  in_channels, 1);
104  TFeat scale = normalizer * n_importance;
105  C.col(out_col) +=
106  (A * (scale * B)).template cast<TOut>();
107  }
108 
109  } // out_idx
110 
111  if (out_importance) {
112  for (int i = 0; i < range_length; ++i)
113  C.col(i) *= TOut(out_importance[r.begin() + i]);
114  }
115  });
116 }
117 
164 template <class TFeat, class TOut, class TIndex, class TKernelIndex>
166  TOut* out_features,
167  const std::vector<int>& filter_dims,
168  const TFeat* filter,
169  size_t num_out,
170  const TFeat* out_importance,
171  size_t num_inp,
172  const TFeat* inp_features,
173  const TFeat* inp_neighbors_importance_sum,
174  const int64_t* inp_neighbors_row_splits,
175  const TIndex* neighbor_index,
176  const TKernelIndex* neighbors_kernel_index,
177  const TFeat* neighbor_importance,
178  const int64_t* neighbors_row_splits,
179  bool normalize) {
180 #define FN_PARAMETERS \
181  out_features, filter_dims, filter, num_out, out_importance, num_inp, \
182  inp_features, inp_neighbors_importance_sum, \
183  inp_neighbors_row_splits, neighbor_index, neighbors_kernel_index, \
184  neighbor_importance, neighbors_row_splits
185 
186 #define CALL_TEMPLATE(NORMALIZE) \
187  if (NORMALIZE == normalize) \
188  _SparseConvTransposeComputeFeaturesCPU<TFeat, TOut, TIndex, \
189  TKernelIndex, NORMALIZE>( \
190  FN_PARAMETERS);
191 
192 #define CALL_TEMPLATE2 \
193  CALL_TEMPLATE(true) \
194  CALL_TEMPLATE(false)
195 
197 
198 #undef CALL_TEMPLATE
199 #undef CALL_TEMPLATE2
200 
201 #undef FN_PARAMETERS
202 }
203 
204 } // namespace impl
205 } // namespace ml
206 } // namespace cloudViewer
#define CALL_TEMPLATE2
__host__ __device__ float2 normalize(float2 v)
Definition: cutil_math.h:1179
void SparseConvTransposeComputeFeaturesCPU(TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, 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 *neighbor_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbor_importance, const int64_t *neighbors_row_splits, bool normalize)
void _SparseConvTransposeComputeFeaturesCPU(TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, 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 *neighbor_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbor_importance, const int64_t *neighbors_row_splits)
Generic file read and write utility for python interface.