kd_node.h

Go to the documentation of this file.

/*
 * Copyright (c) 2010-2016, NVIDIA Corporation
 * All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *   * Redistributions of source code must retain the above copyright
 *     notice, this list of conditions and the following disclaimer.
 *   * Redistributions in binary form must reproduce the above copyright
 *     notice, this list of conditions and the following disclaimer in the
 *     documentation and/or other materials provided with the distribution.
 *   * Neither the name of NVIDIA Corporation nor the
 *     names of its contributors may be used to endorse or promote products
 *     derived from this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL <COPYRIGHT HOLDER> BE LIABLE FOR ANY
 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 */

#pragma once

#include <cugar/basic/types.h>

namespace cugar {


struct Kd_node
{
    static const uint32 kInvalid = uint32(-1);

    CUGAR_HOST_DEVICE Kd_node() {}

    CUGAR_HOST_DEVICE Kd_node(uint32 index) :
        m_packed_info( 7u | (index << 3) ) {}

    CUGAR_HOST_DEVICE Kd_node(const uint32 split_dim, const float split_plane, uint32 index) :
        m_packed_info( split_dim | (index << 3) ),
        m_split_plane( split_plane ) {}

    CUGAR_HOST_DEVICE Kd_node(const uint32 packed_info, const float split_plane) :
        m_packed_info( packed_info ),
        m_split_plane(split_plane) {}
    
    CUGAR_HOST_DEVICE uint32 is_leaf() const
    {
        return (m_packed_info & 7u) == 7u;
    }
    CUGAR_HOST_DEVICE uint32 get_child_offset() const
    {
        return m_packed_info >> 3u;
    }
    CUGAR_HOST_DEVICE uint32 get_leaf_index() const
    {
        return m_packed_info >> 3u;
    }
    CUGAR_HOST_DEVICE uint32 get_child(const uint32 i) const
    {
        return get_child_offset() + i;
    }
    CUGAR_HOST_DEVICE bool has_child(const uint32 i) const
    {
        return is_leaf() ? false : true;
    }
    CUGAR_HOST_DEVICE uint32 get_left() const
    {
        return get_child_offset();
    }
    CUGAR_HOST_DEVICE uint32 get_right() const
    {
        return get_child_offset() + 1u;
    }

    CUGAR_HOST_DEVICE uint32 get_split_dim() const { return (m_packed_info & 7u); }

    CUGAR_HOST_DEVICE float get_split_plane() const { return m_split_plane; }

    CUGAR_HOST_DEVICE
    static Kd_node load(const Kd_node* node)
    {
    #if defined(CUGAR_DEVICE_COMPILATION)
        const uint2 u = *reinterpret_cast<const uint2*>(node);
        return Kd_node(u.x,__uint_as_float(u.y));
    #else
        return *node;
    #endif
    }

    CUGAR_HOST_DEVICE
    static Kd_node load_ldg(const Kd_node* node)
    {
    #if defined(CUGAR_DEVICE_COMPILATION)
        const uint2 u = __ldg(reinterpret_cast<const uint2*>(node));
        return Kd_node(u.x,__uint_as_float(u.y));
    #else
        return *node;
    #endif
    }

    uint32 m_packed_info;
    float  m_split_plane;
};

} // namespace cugar