VertexHashTable.h

#pragma once

// C++ libs
#include <memory>

// CUDA stuff
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

// Project files
#include "helper_cuda.h"

namespace p_mc {
    /// <summary>
    /// Hash table to construct a shared vertex triangle mesh using the
    /// standard Marching Cubes algorithms. Triangles are also constructed
    /// for rendering purposes.
    /// </summary>
    struct VertexHashTable {
        using uint = unsigned int;
        /// size of buffers
        int t_size{ 0 };
        /// key buffer
        int* key{ nullptr };
        std::shared_ptr<int> key_{ nullptr };
        /// index buffer
        int* index{ nullptr };
        std::shared_ptr<int> index_{ nullptr };
        /// constructors
        __host__ VertexHashTable() {}
        __host__ VertexHashTable(const int sz) : t_size{ sz }
        {
            cudaMalloc(&key, sz * sizeof(int));
            cudaMalloc(&index, sz * sizeof(int));
            cudaCheckError();
            key_.reset(key, cudaFree);
            index_.reset(index, cudaFree);
        }
        /// <summary>
        /// Destructor
        /// </summary>
        /// <returns></returns>
        __host__ ~VertexHashTable()
        {
            t_size = 0;
            key = nullptr;
            key_.reset();
            index = nullptr;
            index_.reset();
        }
        /// get size of buffers
        __host__ __device__ int size() { return t_size; }
        /// set default values
        __device__ void init(const int pos)
        {
            key[pos] = EMPTY_BUCKET_32;
            index[pos] = INVALID_INDEX;
        }
        /// simple hash function
        __device__ int hash(const int k)
        {
            return ((3ll * static_cast<unsigned long long>(k)) % 300000007ll) % t_size;
        }
        /// a bit mix hash function
        __device__ uint hash_function(uint key)
        {
            key = (key ^ 61) ^ (key >> 16);
            key = key + (key << 3);
            key = key ^ (key >> 4);
            key = key * 0x27d4eb2d;
            key = key ^ (key >> 15);
            return key;
        }
        /// <summary>
        /// Add vertex into the hash table
        /// </summary>
        /// <param name="k">key</param>
        /// <param name="addr">buffer to keep address of bucket, where vertex was added</param>
        /// <param name="pos">position in vertex buffer</param>
        /// <returns></returns>
        __device__ bool addVertex(const int k, int* addr, const int pos)
        {
            const int start_address = hash(k);
            int h = start_address;
            int e = 1;
            for (int loop = 0; loop < 128; loop++) {
                const int old = atomicCAS(&key[h], EMPTY_BUCKET_32, k);
                //if (old == EMPTY_BUCKET_32 || old == k)
                if (old == EMPTY_BUCKET_32)
                {
                    // vertex has to be created, return bucket address in field
                    addr[pos] = h;
                    return false;
                }
                else if (old == k)
                {
                    // vertex already exists, return bucket address for this key
                    addr[pos] = h;
                    return true;
                }
                else
                {
                    // step with quadratic probing
                    h = (h + e * e) % t_size;
                    e = e + 1;
                    if (h == start_address) {
                        addr[pos] = -1;
                        return false;
                    }
                }
            }
            addr[pos] = -1;
            return false; // something went wrong
        }
        /// access vertex index
        __device__ int v(const int pos) { return index[pos]; }
        /// <summary>
        /// Set vertex index
        /// </summary>
        /// <param name="pos">bucket address</param>
        /// <param name="val">index to be set</param>
        /// <returns></returns>
        __device__ void set(const int pos, const int val) { index[pos] = val; }
    };
} // namespace p_mc