<commit_before>/*
 *  (C) copyright  2011, <NAME>, (U.Girona/ViRVIG, Spain & INRIA/ALICE, France)
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

template <typename T>
struct copy_tex_value
{
  static const libhu::U32 KEY_TYPE_BITS                 = 32u;
  static const libhu::U32 KEY_TYPE_MASK                 = libhu::U32( libhu::U64((1ull) << KEY_TYPE_BITS) - 1ul );
  static const libhu::U32 ATTACH_ID_TYPE_BITS           = 28u;
  static const libhu::U32 ATTACH_ID_TYPE_MASK           = libhu::U32( libhu::U64((1ull) << ATTACH_ID_TYPE_BITS) - 1ul );
  static const libhu::U32 PACKED_UNDEFINED_KEY          = libhu::U32( libhu::U64((1ull) << KEY_TYPE_BITS) - 1ul );

  inline __host__ __device__
    copy_tex_value(libhu::UPTR _keysUPtr, libhu::UPTR _valuesUPtr, libhu::UPTR _refValuesUPtr, libhu::U32 _W, libhu::U32 _H) : keysUPtr(_keysUPtr), valuesUPtr(_valuesUPtr), refValuesUPtr(_refValuesUPtr), W(_W), H(_H) {} 

  //inline __host__ __device__ __forceinline__ libhu::U32 GET_KEY_POS(libhu::U64 k) { return ((k) & KEY_TYPE_MASK); }
  //inline __host__ __device__ __forceinline__ libhu::U32 GET_KEY_ATTACH_ID(libhu::U64 k) { return ((k) >> (KEY_TYPE_BITS)) & ATTACH_ID_TYPE_MASK; }

  #define GET_KEY_POS(k)       ((k) & KEY_TYPE_MASK)
  #define GET_KEY_ATTACH_ID(k) (((k) >> (KEY_TYPE_BITS)) & ATTACH_ID_TYPE_MASK)



  inline __host__ __device__
  void operator()(T& t)
  {
    libhu::U32* valuesPtr = (libhu::U32*)valuesUPtr;
    libhu::U32* refValuesPtr = (libhu::U32*)refValuesUPtr;
    if (t != PACKED_UNDEFINED_KEY)
    {
      //valuesPtr[GET_KEY_POS(t)] = GET_KEY_ATTACH_ID(GET_KEY_POS(t));
      valuesPtr[GET_KEY_POS(t)] = refValuesPtr[GET_KEY_POS(t)];
    }
  }

  libhu::UPTR keysUPtr;
  libhu::UPTR valuesUPtr;
  libhu::UPTR refValuesUPtr;
  libhu::U32 W;
  libhu::U32 H;
};

int testRobinHoodHash(ConfigParams& cfg)
{
  libhu::U32 num_keys     = cfg.num_keys;
  libhu::U32 num_extra    = cfg.num_extra;
  libhu::U32 num_queries  = (cfg.rate_non_valid_keys == 0) ? cfg.num_keys : cfg.num_extra;
  libhu::F32 dens         = cfg.dens;
  libhu::U32 seed         = cfg.seed;
  bool sorted_access      = cfg.sorted_access;
      
  cudaEvent_t start  = 0u;
  cudaEvent_t end    = 0u;
  libhu::F32 coh_build_time;
  libhu::F32 coh_access_time;
  libhu::F32 rand_build_time;
  libhu::F32 rand_access_time;
  
  CUdevice     device;
  bool cvsOutputMode = true;
  
  cudaEventCreate(&start);
  cudaEventCreate(&end);
  
  libhu::U32 hash_table_size = ceilf(num_keys / dens);
  hash_table_size = ((hash_table_size & 1u) == 0u) ? hash_table_size + 1u : hash_table_size;
  
  std::cerr << "mode                             : " << ((ENABLE_KEY_DATA_MODE == 1) ? "key_value (64b)" : "key (32b)") << std::endl;
  std::cerr << "num_keys                         : " << num_keys << std::endl;
  std::cerr << "num_extra                        : " << num_extra << std::endl;
  std::cerr << "num_queries                      : " << num_queries << std::endl;
  std::cerr << "rate_non_valid_keys              : " << cfg.rate_non_valid_keys << std::endl;
  std::cerr << "hash_table_size                  : " << hash_table_size << std::endl;  
  std::cerr << "density                          : " << dens << std::endl;
  std::cerr << "seed                             : " << seed << std::endl;
  std::cerr << "sorted                           : " << sorted_access << std::endl;
  
  typedef libhu::U32 T_KEY;
  typedef libhu::U32 T_VALUE;
  typedef thrust::device_vector<T_KEY>::iterator T_KEY_IT;
  typedef thrust::device_vector<T_VALUE>::iterator T_VALUE_IT;

#if (!ENABLE_KEY_DATA_MODE)
  typedef libhu::U32 T_HASH_TABLE;
  typedef thrust::device_vector<T_HASH_TABLE>::iterator T_HASH_TABLE_IT;
  typedef key_coh_hash_functor  T_COH_FUNCTOR;
  typedef key_rand_hash_functor T_RAND_FUNCTOR;
  key_rand_hash_functor   krf = key_rand_hash_functor();
  key_coh_hash_functor    kcf = key_coh_hash_functor();
  libhu::U32    key_max_value = krf.KEY_TYPE_RANGE;
  libhu::U32      key_max_age = krf.KEY_TYPE_MAX_AGE;
  thrust::device_vector<T_KEY>        values(num_extra);
  thrust::device_vector<T_HASH_TABLE> hash_table(hash_table_size);
#elif (ENABLE_KEY_DATA_MODE)
  typedef libhu::U64 T_HASH_TABLE;
  typedef thrust::device_vector<T_HASH_TABLE>::iterator T_HASH_TABLE_IT;
  typedef key_value_coh_hash_functor  T_COH_FUNCTOR;
  typedef key_value_rand_hash_functor T_RAND_FUNCTOR;
  key_value_rand_hash_functor krf = key_value_rand_hash_functor();
  key_value_coh_hash_functor  kcf = key_value_coh_hash_functor();
  libhu::U32        key_max_value = krf.KEY_TYPE_RANGE;
  libhu::U32          key_max_age = krf.KEY_TYPE_MAX_AGE;
  thrust::device_vector<T_VALUE>      values(num_extra);
  thrust::device_vector<T_HASH_TABLE> hash_table(hash_table_size);
#endif

  thrust::device_vector<T_KEY> keys(num_extra);
  thrust::fill(keys.begin(), keys.end(), kcf.UNDEFINED_KEY);

  if (cfg.rand_num_mode)
  {

#if (ENABLE_LIBHU_GPU_RANDOM_GENERATOR)
  
#if (ENABLE_2D_COORDS)
    libhu::key_2d<libhu::U32, kcf.KEY_TYPE_BITS> key_2d_converter;
    key_max_value = (key_2d_converter.w * key_2d_converter.h) - 1u;
    //libhu::key_2d<libhu::U32, kcf.KEY_TYPE_BITS> key_2d_converter(cfg.u2D_w, cfg.u2D_h);
    //key_max_value = (key_2d_converter.w * key_2d_converter.h) - 1u;
    if (key_max_value < cfg.num_extra) 
    {
      std::cerr << "Error: num_extra > universe size" << std::endl;
      exit(0);
    }
    bool unique = libhu::generate_unique_random(keys.begin(), keys.end(), 1u, key_max_value, seed);
    std::cerr << "2D universe (" << key_2d_converter.w << "," << key_2d_converter.h << ")" << std::endl;
#elif (ENABLE_3D_COORDS)
    libhu::key_3d<libhu::U32, kcf.KEY_TYPE_BITS> key_3d_converter;  <commit_msg>Checking problem with class static const<commit_after>/*
 *  (C) copyright  2011, <NAME>, (U.Girona/ViRVIG, Spain & INRIA/ALICE, France)
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

template <typename T>
struct copy_tex_value
{
  static const libhu::U32 KEY_TYPE_BITS                 = 32u;
  static const libhu::U32 KEY_TYPE_MASK                 = libhu::U32( libhu::U64((1ull) << KEY_TYPE_BITS) - 1ul );
  static const libhu::U32 ATTACH_ID_TYPE_BITS           = 28u;
  static const libhu::U32 ATTACH_ID_TYPE_MASK           = libhu::U32( libhu::U64((1ull) << ATTACH_ID_TYPE_BITS) - 1ul );
  static const libhu::U32 PACKED_UNDEFINED_KEY          = libhu::U32( libhu::U64((1ull) << KEY_TYPE_BITS) - 1ul );

  inline __host__ __device__
    copy_tex_value(libhu::UPTR _keysUPtr, libhu::UPTR _valuesUPtr, libhu::UPTR _refValuesUPtr, libhu::U32 _W, libhu::U32 _H) : keysUPtr(_keysUPtr), valuesUPtr(_valuesUPtr), refValuesUPtr(_refValuesUPtr), W(_W), H(_H) {} 

  //inline __host__ __device__ __forceinline__ libhu::U32 GET_KEY_POS(libhu::U64 k) { return ((k) & KEY_TYPE_MASK); }
  //inline __host__ __device__ __forceinline__ libhu::U32 GET_KEY_ATTACH_ID(libhu::U64 k) { return ((k) >> (KEY_TYPE_BITS)) & ATTACH_ID_TYPE_MASK; }

  #define GET_KEY_POS(k)       ((k) & KEY_TYPE_MASK)
  #define GET_KEY_ATTACH_ID(k) (((k) >> (KEY_TYPE_BITS)) & ATTACH_ID_TYPE_MASK)



  inline __host__ __device__
  void operator()(T& t)
  {
    libhu::U32* valuesPtr = (libhu::U32*)valuesUPtr;
    libhu::U32* refValuesPtr = (libhu::U32*)refValuesUPtr;
    if (t != PACKED_UNDEFINED_KEY)
    {
      //valuesPtr[GET_KEY_POS(t)] = GET_KEY_ATTACH_ID(GET_KEY_POS(t));
      valuesPtr[GET_KEY_POS(t)] = refValuesPtr[GET_KEY_POS(t)];
    }
  }

  libhu::UPTR keysUPtr;
  libhu::UPTR valuesUPtr;
  libhu::UPTR refValuesUPtr;
  libhu::U32 W;
  libhu::U32 H;
};

int testRobinHoodHash(ConfigParams& cfg)
{
  libhu::U32 num_keys     = cfg.num_keys;
  libhu::U32 num_extra    = cfg.num_extra;
  libhu::U32 num_queries  = (cfg.rate_non_valid_keys == 0) ? cfg.num_keys : cfg.num_extra;
  libhu::F32 dens         = cfg.dens;
  libhu::U32 seed         = cfg.seed;
  bool sorted_access      = cfg.sorted_access;
      
  cudaEvent_t start  = 0u;
  cudaEvent_t end    = 0u;
  libhu::F32 coh_build_time;
  libhu::F32 coh_access_time;
  libhu::F32 rand_build_time;
  libhu::F32 rand_access_time;
  
  CUdevice     device;
  bool cvsOutputMode = true;
  
  cudaEventCreate(&start);
  cudaEventCreate(&end);
  
  libhu::U32 hash_table_size = ceilf(num_keys / dens);
  hash_table_size = ((hash_table_size & 1u) == 0u) ? hash_table_size + 1u : hash_table_size;
  
  std::cerr << "mode                             : " << ((ENABLE_KEY_DATA_MODE == 1) ? "key_value (64b)" : "key (32b)") << std::endl;
  std::cerr << "num_keys                         : " << num_keys << std::endl;
  std::cerr << "num_extra                        : " << num_extra << std::endl;
  std::cerr << "num_queries                      : " << num_queries << std::endl;
  std::cerr << "rate_non_valid_keys              : " << cfg.rate_non_valid_keys << std::endl;
  std::cerr << "hash_table_size                  : " << hash_table_size << std::endl;  
  std::cerr << "density                          : " << dens << std::endl;
  std::cerr << "seed                             : " << seed << std::endl;
  std::cerr << "sorted                           : " << sorted_access << std::endl;
  
  typedef libhu::U32 T_KEY;
  typedef libhu::U32 T_VALUE;
  typedef thrust::device_vector<T_KEY>::iterator T_KEY_IT;
  typedef thrust::device_vector<T_VALUE>::iterator T_VALUE_IT;

#if (!ENABLE_KEY_DATA_MODE)
  typedef libhu::U32 T_HASH_TABLE;
  typedef thrust::device_vector<T_HASH_TABLE>::iterator T_HASH_TABLE_IT;
  typedef key_coh_hash_functor  T_COH_FUNCTOR;
  typedef key_rand_hash_functor T_RAND_FUNCTOR;
  key_rand_hash_functor   krf = key_rand_hash_functor();
  key_coh_hash_functor    kcf = key_coh_hash_functor();
  libhu::U32    key_max_value = krf.KEY_TYPE_RANGE;
  libhu::U32      key_max_age = krf.KEY_TYPE_MAX_AGE;
  thrust::device_vector<T_KEY>        values(num_extra);
  thrust::device_vector<T_HASH_TABLE> hash_table(hash_table_size);
#elif (ENABLE_KEY_DATA_MODE)
  typedef libhu::U64 T_HASH_TABLE;
  typedef thrust::device_vector<T_HASH_TABLE>::iterator T_HASH_TABLE_IT;
  typedef key_value_coh_hash_functor  T_COH_FUNCTOR;
  typedef key_value_rand_hash_functor T_RAND_FUNCTOR;
  key_value_rand_hash_functor krf = key_value_rand_hash_functor();
  key_value_coh_hash_functor  kcf = key_value_coh_hash_functor();
  libhu::U32        key_max_value = krf.KEY_TYPE_RANGE;
  libhu::U32          key_max_age = krf.KEY_TYPE_MAX_AGE;
  thrust::device_vector<T_VALUE>      values(num_extra);
  thrust::device_vector<T_HASH_TABLE> hash_table(hash_table_size);
#endif

  thrust::device_vector<T_KEY> keys(num_extra);
  thrust::fill(keys.begin(), keys.end(), T_COH_FUNCTOR::UNDEFINED_KEY);//kcf.UNDEFINED_KEY);

  if (cfg.rand_num_mode)
  {

#if (ENABLE_LIBHU_GPU_RANDOM_GENERATOR)
  
#if (ENABLE_2D_COORDS)
    libhu::key_2d<libhu::U32, kcf.KEY_TYPE_BITS> key_2d_converter;
    key_max_value = (key_2d_converter.w * key_2d_converter.h) - 1u;
    //libhu::key_2d<libhu::U32, kcf.KEY_TYPE_BITS> key_2d_converter(cfg.u2D_w, cfg.u2D_h);
    //key_max_value = (key_2d_converter.w * key_2d_converter.h) - 1u;
    if (key_max_value < cfg.num_extra) 
    {
      std::cerr << "Error: num_extra > universe size" << std::endl;
      exit(0);
    }
    bool unique = libhu::generate_unique_random(keys.begin(), keys.end(), 1u, key_max_value, seed);
    std::cerr << "2D universe (" << key_2d_converter.w << "," << key_2d_converter.h << ")" << std::endl;
#elif (ENABLE_3D_COORDS)
    libhu::key_3d<libhu::U32, kcf.KEY_TYPE_BITS> key_3d_converter;  
{'fraction_non_alphanumeric': 0.12516985053153226, 'fraction_numerical': 0.02206058668371833, 'mean_word_length': 3.782874617737003, 'pattern_counts': {'":': 0, '<': 185, '<?xml version=': 0, '>': 41, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 6, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '1989144', 'n_tokens_mistral': 4945, 'n_tokens_neox': 4761, 'n_words': 859}