BlockCopyDispatch.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>
 
#ifdef SYCL_LANGUAGE_VERSION
#include <sycl/sycl.hpp>
#endif
 
namespace open3d {
namespace core {
 
// Vectorized trivial-object copy block sizes (bytes). Used for: HashMap SoA
// value buffers (VoxelBlockGrid blocks, etc.), Dtype::Object advanced index
// get/set, and non-contiguous SYCL tensor copy of object dtypes.
inline constexpr int64_t kBlockCopyDivisors[] = {64, 16, 12, 4, 1};
 
inline int64_t GetLargestAlignedObjectBlockSize(int64_t object_byte_size) {
    for (int64_t divisor : kBlockCopyDivisors) {
        if (object_byte_size % divisor == 0) {
            return divisor;
        }
    }
    return 1;
}
 
}  // namespace core
}  // namespace open3d
 
#ifdef __CUDACC__
#include <cuda_runtime.h>
 
// Reinterpret hash maps' void* value arrays as CUDA primitive type arrays to
// avoid slow memcpy or byte-by-byte copy in kernels. Not used on CPU (memcpy is
// fast enough). BlockCopy64 is at namespace scope because nvcc disallows
// types with no linkage as template arguments for __global__ instantiations.
struct BlockCopy64 {
    int4 v[4];
};
 
#define DISPATCH_DIVISOR_SIZE_TO_BLOCK_T(DIVISOR, ...)                     \
    [&] {                                                                  \
        if (DIVISOR == 64) {                                               \
            using block_t = BlockCopy64;                                   \
            static_assert(sizeof(block_t) == 64, "block_t size mismatch"); \
            return __VA_ARGS__();                                          \
        } else if (DIVISOR == 16) {                                        \
            using block_t = int4;                                          \
            static_assert(sizeof(block_t) == 16, "block_t size mismatch"); \
            return __VA_ARGS__();                                          \
        } else if (DIVISOR == 12) {                                        \
            using block_t = int3;                                          \
            static_assert(sizeof(block_t) == 12, "block_t size mismatch"); \
            return __VA_ARGS__();                                          \
        } else if (DIVISOR == 4) {                                         \
            using block_t = int;                                           \
            static_assert(sizeof(block_t) == 4, "block_t size mismatch");  \
            return __VA_ARGS__();                                          \
        } else {                                                           \
            using block_t = uint8_t;                                       \
            static_assert(sizeof(block_t) == 1, "block_t size mismatch");  \
            return __VA_ARGS__();                                          \
        }                                                                  \
    }()
#endif  // __CUDACC__
 
#ifdef SYCL_LANGUAGE_VERSION
 
// sycl::vec<uint32_t, 3> is NOT 12 bytes: the SYCL spec requires 3-element
// vectors to have the same size/alignment as the 4-element vector (16 bytes,
// 16-byte aligned), so it cannot be used for the 12-byte divisor. Use a
// plain POD struct instead, which packs to exactly 12 bytes.
struct BlockCopy12 {
    uint32_t v[3];
};
 
#define DISPATCH_DIVISOR_SIZE_TO_BLOCK_T_SYCL(DIVISOR, ...)                \
    [&] {                                                                  \
        if (DIVISOR == 64) {                                               \
            using block_t = sycl::vec<uint32_t, 16>;                       \
            static_assert(sizeof(block_t) == 64, "block_t size mismatch"); \
            return __VA_ARGS__();                                          \
        } else if (DIVISOR == 16) {                                        \
            using block_t = sycl::vec<uint32_t, 4>;                        \
            static_assert(sizeof(block_t) == 16, "block_t size mismatch"); \
            return __VA_ARGS__();                                          \
        } else if (DIVISOR == 12) {                                        \
            using block_t = BlockCopy12;                                   \
            static_assert(sizeof(block_t) == 12, "block_t size mismatch"); \
            return __VA_ARGS__();                                          \
        } else if (DIVISOR == 4) {                                         \
            using block_t = uint32_t;                                      \
            static_assert(sizeof(block_t) == 4, "block_t size mismatch");  \
            return __VA_ARGS__();                                          \
        } else {                                                           \
            using block_t = uint8_t;                                       \
            static_assert(sizeof(block_t) == 1, "block_t size mismatch");  \
            return __VA_ARGS__();                                          \
        }                                                                  \
    }()
 
#endif  // SYCL_LANGUAGE_VERSION