ParallelFor.h

Go to the documentation of this file.

// ----------------------------------------------------------------------------
// -                        Open3D: www.open3d.org                            -
// ----------------------------------------------------------------------------
// Copyright (c) 2018-2026 www.open3d.org
// SPDX-License-Identifier: MIT
// ----------------------------------------------------------------------------
 
#pragma once
 
#include <cstdint>
#include <type_traits>
 
#include "open3d/core/Device.h"
#include "open3d/utility/Logging.h"
#include "open3d/utility/Overload.h"
#include "open3d/utility/Parallel.h"
#include "open3d/utility/Preprocessor.h"
 
#ifdef __CUDACC__
#include <cuda.h>
#include <cuda_runtime.h>
 
#include "open3d/core/CUDAUtils.h"
#else
#include <tbb/parallel_for.h>
#endif
 
#if defined(SYCL_LANGUAGE_VERSION)
#include "open3d/core/SYCLUtils.h"
#endif
 
namespace open3d {
namespace core {
 
#ifdef __CUDACC__
 
static constexpr int64_t OPEN3D_PARFOR_BLOCK = 128;
static constexpr int64_t OPEN3D_PARFOR_THREAD = 4;
 
template <int64_t block_size, int64_t thread_size, typename func_t>
__global__ void ElementWiseKernel_(int64_t n, func_t f) {
    int64_t items_per_block = block_size * thread_size;
    int64_t idx = blockIdx.x * items_per_block + threadIdx.x;
#pragma unroll
    for (int64_t i = 0; i < thread_size; ++i) {
        if (idx < n) {
            f(idx);
            idx += block_size;
        }
    }
}
 
template <typename func_t>
void ParallelForCUDA_(const Device& device, int64_t n, const func_t& func) {
    if (device.GetType() != Device::DeviceType::CUDA) {
        utility::LogError("ParallelFor for CUDA cannot run on device {}.",
                          device.ToString());
    }
    if (n == 0) {
        return;
    }
 
    CUDAScopedDevice scoped_device(device);
    int64_t items_per_block = OPEN3D_PARFOR_BLOCK * OPEN3D_PARFOR_THREAD;
    int64_t grid_size = (n + items_per_block - 1) / items_per_block;
 
    ElementWiseKernel_<OPEN3D_PARFOR_BLOCK, OPEN3D_PARFOR_THREAD>
            <<<grid_size, OPEN3D_PARFOR_BLOCK, 0, core::cuda::GetStream()>>>(
                    n, func);
    OPEN3D_GET_LAST_CUDA_ERROR("ParallelFor failed.");
}
 
#else
 
template <typename func_t>
void ParallelForCPU_(const Device& device,
                     int64_t n,
                     int64_t grain_size,
                     const func_t& func) {
    if (!device.IsCPU()) {
        utility::LogError("ParallelFor for CPU cannot run on device {}.",
                          device.ToString());
    }
    if (n == 0) {
        return;
    }
 
    tbb::parallel_for(tbb::blocked_range<int64_t>(0, n, grain_size),
                      [&func](const tbb::blocked_range<int64_t>& range) {
                          for (int64_t i = range.begin(); i < range.end();
                               ++i) {
                              func(i);
                          }
                      });
}
 
template <typename func_t>
void ParallelForCPU_(const Device& device, int64_t n, const func_t& func) {
    ParallelForCPU_(device, n,
                    static_cast<int64_t>(utility::DefaultGrainSizeTBB2D()),
                    func);
}
 
#endif
 
#if defined(SYCL_LANGUAGE_VERSION)
 
template <typename func_t>
sycl::event ParallelForSYCLImpl_(sycl::queue& queue,
                                 int64_t n,
                                 const func_t& func,
                                 const std::vector<sycl::event>& deps) {
    if (n == 0) {
        return sycl::event();
    }
    size_t wg = core::sy::PreferredWorkGroupSize(queue.get_device());
    const size_t global_size = ((static_cast<size_t>(n) + wg - 1) / wg) * wg;
    sycl::nd_range<1> nd_range{sycl::range<1>(global_size), sycl::range<1>(wg)};
    return queue.parallel_for(nd_range, deps, [=](sycl::nd_item<1> item) {
        int64_t i = item.get_global_id(0);
        if (i < n) {
            func(i);
        }
    });
}
 
template <typename func_t>
void ParallelForSYCL_(const Device& device, int64_t n, const func_t& func) {
    if (!device.IsSYCL()) {
        utility::LogError("ParallelFor for SYCL cannot run on device {}.",
                          device.ToString());
    }
    auto queue = core::sy::GetQueue(device);
    ParallelForSYCLImpl_(queue, n, func, {}).wait_and_throw();
}
 
template <typename func_t>
sycl::event ParallelForSYCL_(const Device& device,
                             int64_t n,
                             const func_t& func,
                             const std::vector<sycl::event>& deps) {
    if (!device.IsSYCL()) {
        utility::LogError("ParallelFor for SYCL cannot run on device {}.",
                          device.ToString());
    }
    auto queue = core::sy::GetQueue(device);
    return ParallelForSYCLImpl_(queue, n, func, deps);
}
 
#endif
 
template <typename func_t>
void ParallelFor(const Device& device, int64_t n, const func_t& func) {
#ifdef __CUDACC__
    ParallelForCUDA_(device, n, func);
#elif defined(SYCL_LANGUAGE_VERSION)
    if (device.IsSYCL()) {
        ParallelForSYCL_(device, n, func);
    } else {
        ParallelForCPU_(device, n, func);
    }
#else
    ParallelForCPU_(device, n, func);
#endif
}
 
#if defined(SYCL_LANGUAGE_VERSION)
template <typename func_t>
sycl::event ParallelFor(const Device& device,
                        int64_t n,
                        const func_t& func,
                        std::initializer_list<sycl::event> deps) {
    return ParallelForSYCL_(device, n, func, std::vector<sycl::event>(deps));
}
 
template <typename func_t>
void ParallelFor(sycl::queue& queue, int64_t n, const func_t& func) {
    ParallelForSYCLImpl_(queue, n, func, {}).wait_and_throw();
}
 
template <typename func_t>
sycl::event ParallelFor(sycl::queue& queue,
                        int64_t n,
                        const func_t& func,
                        const std::vector<sycl::event>& deps) {
    return ParallelForSYCLImpl_(queue, n, func, deps);
}
 
template <typename func_t>
sycl::event ParallelFor(sycl::queue& queue,
                        int64_t n,
                        const func_t& func,
                        std::initializer_list<sycl::event> deps) {
    return ParallelFor(queue, n, func, std::vector<sycl::event>(deps));
}
#endif
 
template <typename vec_func_t, typename func_t>
void ParallelFor(const Device& device,
                 int64_t n,
                 const func_t& func,
                 const vec_func_t& vec_func) {
#ifdef BUILD_ISPC_MODULE
 
#ifdef __CUDACC__
    ParallelForCUDA_(device, n, func);
#elif defined(SYCL_LANGUAGE_VERSION)
    if (device.IsSYCL()) {
        ParallelForSYCL_(device, n, func);
    } else {
        // The work is already blocked into num_threads chunks here, so the
        // grain size must be 1: a larger grain would exceed the range and
        // make TBB run every chunk on a single thread.
        int num_threads = utility::EstimateMaxThreads();
        ParallelForCPU_(device, num_threads, /*grain_size=*/1, [&](int64_t i) {
            int64_t start = n * i / num_threads;
            int64_t end = std::min<int64_t>(n * (i + 1) / num_threads, n);
            vec_func(start, end);
        });
    }
#else
    // See the comment above: the range is already one item per thread.
    int num_threads = utility::EstimateMaxThreads();
    ParallelForCPU_(device, num_threads, /*grain_size=*/1, [&](int64_t i) {
        int64_t start = n * i / num_threads;
        int64_t end = std::min<int64_t>(n * (i + 1) / num_threads, n);
        vec_func(start, end);
    });
#endif
 
#else
 
#ifdef __CUDACC__
    ParallelForCUDA_(device, n, func);
#elif defined(SYCL_LANGUAGE_VERSION)
    if (device.IsSYCL()) {
        ParallelForSYCL_(device, n, func);
    } else {
        ParallelForCPU_(device, n, func);
    }
#else
    ParallelForCPU_(device, n, func);
#endif
 
#endif
}
 
#ifdef BUILD_ISPC_MODULE
 
// Internal helper macro.
#define OPEN3D_CALL_ISPC_KERNEL_(ISPCKernel, start, end, ...) \
    using namespace ispc;                                     \
    ISPCKernel(start, end, __VA_ARGS__);
 
#else
 
// Internal helper macro.
#define OPEN3D_CALL_ISPC_KERNEL_(ISPCKernel, start, end, ...)            \
    utility::LogError(                                                   \
            "ISPC module disabled. Unable to call vectorized kernel {}", \
            OPEN3D_STRINGIFY(ISPCKernel));
 
#endif
 
#define OPEN3D_OVERLOADED_LAMBDA_(T, ISPCKernel, ...)                       \
    [&](T, int64_t start, int64_t end) {                                    \
        OPEN3D_CALL_ISPC_KERNEL_(                                           \
                OPEN3D_CONCAT(ISPCKernel, OPEN3D_CONCAT(_, T)), start, end, \
                __VA_ARGS__);                                               \
    }
 
#define OPEN3D_VECTORIZED(ISPCKernel, ...)                             \
    [&](int64_t start, int64_t end) {                                  \
        OPEN3D_CALL_ISPC_KERNEL_(ISPCKernel, start, end, __VA_ARGS__); \
    }
 
#define OPEN3D_TEMPLATE_VECTORIZED(T, ISPCKernel, ...)                        \
    [&](int64_t start, int64_t end) {                                         \
        static_assert(std::is_arithmetic<T>::value,                           \
                      "Data type is not an arithmetic type");                 \
        utility::Overload(                                                    \
                OPEN3D_OVERLOADED_LAMBDA_(bool, ISPCKernel, __VA_ARGS__),     \
                OPEN3D_OVERLOADED_LAMBDA_(uint8_t, ISPCKernel, __VA_ARGS__),  \
                OPEN3D_OVERLOADED_LAMBDA_(int8_t, ISPCKernel, __VA_ARGS__),   \
                OPEN3D_OVERLOADED_LAMBDA_(uint16_t, ISPCKernel, __VA_ARGS__), \
                OPEN3D_OVERLOADED_LAMBDA_(int16_t, ISPCKernel, __VA_ARGS__),  \
                OPEN3D_OVERLOADED_LAMBDA_(uint32_t, ISPCKernel, __VA_ARGS__), \
                OPEN3D_OVERLOADED_LAMBDA_(int32_t, ISPCKernel, __VA_ARGS__),  \
                OPEN3D_OVERLOADED_LAMBDA_(uint64_t, ISPCKernel, __VA_ARGS__), \
                OPEN3D_OVERLOADED_LAMBDA_(int64_t, ISPCKernel, __VA_ARGS__),  \
                OPEN3D_OVERLOADED_LAMBDA_(float, ISPCKernel, __VA_ARGS__),    \
                OPEN3D_OVERLOADED_LAMBDA_(double, ISPCKernel, __VA_ARGS__),   \
                [&](auto&& generic, int64_t start, int64_t end) {             \
                    utility::LogError(                                        \
                            "Unsupported data type {} for calling "           \
                            "vectorized kernel {}",                           \
                            typeid(generic).name(),                           \
                            OPEN3D_STRINGIFY(ISPCKernel));                    \
                })(T{}, start, end);                                          \
    }
 
}  // namespace core
}  // namespace open3d