10__device__
size_t get_thread_id() {
11 return static_cast<size_t>(blockIdx.x) *
static_cast<size_t>(blockDim.x) +
12 static_cast<size_t>(threadIdx.x);
15template <
typename T,
size_t I,
size_t N,
typename M,
size_t E,
typename F,
16 typename VT, std::size_t... Is>
17__device__
void compute_Jblock(
18 T *jacobian,
const size_t factor_id,
const size_t vertex_id,
const M *obs,
19 const typename F::ConstraintDataType *constraint_data,
const size_t *ids,
20 const size_t *hessian_ids, VT args, std::index_sequence<Is...>) {
22 constexpr auto vertex_sizes = F::get_vertex_sizes();
24 const M *local_obs = obs + factor_id;
25 const typename F::ConstraintDataType *local_data =
26 constraint_data + factor_id;
28 auto vargs = cuda::std::make_tuple(
29 (*(std::get<Is>(args) + ids[factor_id * N + Is]))...);
31 using DataType =
typename F::ConstraintDataType;
32 using ObsType =
typename F::ObservationType;
34 if constexpr (std::is_empty<ObsType>::value &&
35 std::is_empty<DataType>::value) {
36 F::Traits::template jacobian<T, I>((*cuda::std::get<Is>(vargs))...,
38 }
else if constexpr (std::is_empty<DataType>::value) {
39 F::Traits::template jacobian<T, I>((*cuda::std::get<Is>(vargs))...,
40 *local_obs, jacobian);
41 }
else if constexpr (std::is_empty<ObsType>::value) {
42 F::Traits::template jacobian<T, I>((*cuda::std::get<Is>(vargs))...,
43 *local_data, jacobian);
45 F::Traits::template jacobian<T, I>((*cuda::std::get<Is>(vargs))...,
46 *local_obs, *local_data, jacobian);