Registration Reduction6x6Impl template datatype support.

[reyanshsolis]

Changes:
Adding template dtype support in Reduction6x6Impl.

Issue re-solved [by @stotko ]:
Adding

#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double *address, double val) {
    unsigned long long int *address_as_ull = (unsigned long long int *)address;
    unsigned long long int old = *address_as_ull, assumed;

    do {
        assumed = old;
        old = atomicCAS(
                address_as_ull, assumed,
                __double_as_longlong(val + __longlong_as_double(assumed)));

        // Note: uses integer comparison to avoid hang in case of NaN (since NaN
        // != NaN)
    } while (assumed != old);

    return __longlong_as_double(old);
}
#endif

throws error

.../Open3D/cpp/open3d/t/geometry/kernel/GeometryMacros.h(35): error: function "atomicAdd(double *, double)" has already been defined

Not adding the same, throws an error in some of the CIs, that atomicAdd(double *, double) is not defined.
[Refer to the CI failure in PR #3564 ]

D:\a\Open3D\Open3D\cpp\open3d/t/pipelines/kernel/Reduction6x6Impl.cuh(180): error : no instance of overloaded function "atomicAdd" matches the argument list [C:\Open3D\build\cpp\open3d\t\pipelines\kernel\tpipelines_kernel.vcxproj]
              argument types are: (double *, volatile double)
            detected during:
              instantiation of "void open3d::t::pipelines::kernel::ReduceSum6x6LinearSystem<T,BLOCK_SIZE>(int, __nv_bool, const T *, volatile T *, volatile T *, volatile T *, T *) [with T=double, BLOCK_SIZE=256ULL]" 

---

This change is 

[update-docs]

Thanks for submitting this pull request! The maintainers of this repository would appreciate if you could update the CHANGELOG.md based on your changes.

[reyanshsolis]

Thanks, @stotko for the help. I have changed the topic of this pr for template datatype support in the Reduction6x6Impl code, which is used later in my other PR for adding F64 support in registration. Could you please review for the same ?

[stotko]

Reviewed 1 of 2 files at r1, 1 of 1 files at r2.
Reviewable status: all files reviewed, 1 unresolved discussion (waiting on @reyanshsolis)

---

cpp/open3d/t/pipelines/kernel/Reduction6x6Impl.cuh, line 168 at r2 (raw file):

        if (tid == 0) {
            OPEN3D_ATOMIC_ADD(&global_sum[i + 0], local_sum0[0]);

nit: It's fine to keep the old non-portable atomicAdd calls here (and below) since the function will not work on __host__ at the moment.

[reyanshsolis]

Reviewed 1 of 2 files at r1, 1 of 1 files at r2.
Reviewable status: all files reviewed, 1 unresolved discussion (waiting on @reyanshsolis)

cpp/open3d/t/pipelines/kernel/Reduction6x6Impl.cuh, line 168 at r2 (raw file):

        if (tid == 0) {
            OPEN3D_ATOMIC_ADD(&global_sum[i + 0], local_sum0[0]);

nit: It's fine to keep the old non-portable atomicAdd calls here (and below) since the function will not work on __host__ at the moment.

Done.