Fix a bug in NonZeroCUDA not using the correct memory addresses (#7340)

---------
Co-authored-by: Sameer Sheorey <[email protected]>

Author: TwentyPast4

CHANGELOG.md

@@ -63,6 +63,7 @@
 -   Exposed `get_plotly_fig` and modified `draw_plotly` to return the `Figure` it creates. (PR #7258)
 -   Fix build with librealsense v2.44.0 and upcoming VS 2022 17.13 (PR #7074)
 -   Fix `deprecated-declarations` warnings when compiling code with C++20 standard (PR #7303)
+-   Fix advanced indexing bug with sliced boolean masks on CUDA devices (PR #7340)
 -   Fix logic for adding -allow-unsupported-compiler to nvcc (PR #7337)
 -   Fix linker error "library limit of 65535 objects exceeded" with Ninja generator on MSVC (PR #7335)
 

cpp/open3d/core/kernel/NonZeroCUDA.cu

@@ -69,8 +69,8 @@ Tensor NonZeroCUDA(const Tensor& src) {
     // Get flattened non-zero indices.
     thrust::device_vector<int64_t> non_zero_indices(num_elements);
     DISPATCH_DTYPE_TO_TEMPLATE_WITH_BOOL(src.GetDtype(), [&]() {
-        thrust::device_ptr<const scalar_t> src_ptr(static_cast<const scalar_t*>(
-                src_contiguous.GetBlob()->GetDataPtr()));
+        thrust::device_ptr<const scalar_t> src_ptr(
+                static_cast<const scalar_t*>(src_contiguous.GetDataPtr()));
 
         auto it = thrust::copy_if(index_first, index_last, src_ptr,
                                   non_zero_indices.begin(),