From b6d6b2905259f974da4ce6a4df2f6f1bce487c55 Mon Sep 17 00:00:00 2001
From: davidwendt <dwendt@nvidia.com>
Date: Tue, 25 May 2021 10:00:56 -0400
Subject: [PATCH 1/3] Fix concurrent_unordered_map to prevent accessing padding
 bits in pair_type

---
 cpp/src/hash/concurrent_unordered_map.cuh | 22 +++++++++++-----------
 1 file changed, 11 insertions(+), 11 deletions(-)

diff --git a/cpp/src/hash/concurrent_unordered_map.cuh b/cpp/src/hash/concurrent_unordered_map.cuh
index 2dfd15925d2..f9e9f09aca4 100644
--- a/cpp/src/hash/concurrent_unordered_map.cuh
+++ b/cpp/src/hash/concurrent_unordered_map.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017-2020, NVIDIA CORPORATION.  All rights reserved.
+ * Copyright (c) 2017-2021, NVIDIA CORPORATION.  All rights reserved.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,14 +14,13 @@
  * limitations under the License.
  */
 
-#ifndef CONCURRENT_UNORDERED_MAP_CUH
-#define CONCURRENT_UNORDERED_MAP_CUH
+#pragma once
 
-#include <cudf/detail/nvtx/ranges.hpp>
 #include <hash/hash_allocator.cuh>
 #include <hash/helper_functions.cuh>
 #include <hash/managed.cuh>
 
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/device_atomics.cuh>
 #include <cudf/detail/utilities/hash_functions.cuh>
 #include <cudf/utilities/error.hpp>
@@ -30,7 +29,6 @@
 
 #include <thrust/pair.h>
 
-#include <cassert>
 #include <iostream>
 #include <iterator>
 #include <limits>
@@ -59,11 +57,15 @@ using packed_t = typename packed<sizeof(pair_type)>::type;
  * equal in size to a type where atomicCAS is natively supported, it is more
  * efficient to "pack" the pair and insert it with a single atomicCAS.
  *
- * @note Only integral key and value types may be packed because we use
+ * Only integral key and value types may be packed because we use
  * bitwise equality comparison, which may not be valid for non-integral
  * types.
  *
- * @tparam pair_type The pair type in question
+ * Also, the `key_type` and `value_type` must be the same size otherwise
+ * the packed value will contain padding bits and accessing the packed
+ * value would be undefined.
+ *
+ * @tparam pair_type The pair type that will be packed
  * @return true If the pair type can be packed
  * @return false  If the pair type cannot be packed
  */
@@ -73,7 +75,7 @@ template <typename pair_type,
 constexpr bool is_packable()
 {
   return std::is_integral<key_type>::value and std::is_integral<value_type>::value and
-         not std::is_void<packed_t<pair_type>>::value;
+         not std::is_void<packed_t<pair_type>>::value and sizeof(key_type) == sizeof(value_type);
 }
 
 /**
@@ -263,7 +265,7 @@ class concurrent_unordered_map {
    */
   template <typename pair_type = value_type>
   __device__ std::enable_if_t<is_packable<pair_type>(), insert_result> attempt_insert(
-    value_type* insert_location, value_type const& insert_pair)
+    value_type* const __restrict__ insert_location, value_type const& insert_pair)
   {
     pair_packer<pair_type> const unused{thrust::make_pair(m_unused_key, m_unused_element)};
     pair_packer<pair_type> const new_pair{insert_pair};
@@ -541,5 +543,3 @@ class concurrent_unordered_map {
     CUDA_TRY(cudaGetLastError());
   }
 };
-
-#endif  // CONCURRENT_UNORDERED_MAP_CUH

From ac8730fea7f80a00d67996aba7d8b3f1872d7b9f Mon Sep 17 00:00:00 2001
From: davidwendt <dwendt@nvidia.com>
Date: Tue, 25 May 2021 10:41:26 -0400
Subject: [PATCH 2/3] use std::has_unique_object_representations instead of
 sizeof check

---
 cpp/src/hash/concurrent_unordered_map.cuh | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/cpp/src/hash/concurrent_unordered_map.cuh b/cpp/src/hash/concurrent_unordered_map.cuh
index f9e9f09aca4..8bf36f41ae9 100644
--- a/cpp/src/hash/concurrent_unordered_map.cuh
+++ b/cpp/src/hash/concurrent_unordered_map.cuh
@@ -75,7 +75,8 @@ template <typename pair_type,
 constexpr bool is_packable()
 {
   return std::is_integral<key_type>::value and std::is_integral<value_type>::value and
-         not std::is_void<packed_t<pair_type>>::value and sizeof(key_type) == sizeof(value_type);
+         not std::is_void<packed_t<pair_type>>::value and
+         std::has_unique_object_representations_v<pair_type>;
 }
 
 /**

From e7c668b6379a639c45c71cc3fc1527973f5f9b6f Mon Sep 17 00:00:00 2001
From: davidwendt <dwendt@nvidia.com>
Date: Tue, 25 May 2021 20:06:15 -0400
Subject: [PATCH 3/3] fix detail comment for is_packable

---
 cpp/src/hash/concurrent_unordered_map.cuh | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/cpp/src/hash/concurrent_unordered_map.cuh b/cpp/src/hash/concurrent_unordered_map.cuh
index 8bf36f41ae9..c4a9da9285d 100644
--- a/cpp/src/hash/concurrent_unordered_map.cuh
+++ b/cpp/src/hash/concurrent_unordered_map.cuh
@@ -61,9 +61,8 @@ using packed_t = typename packed<sizeof(pair_type)>::type;
  * bitwise equality comparison, which may not be valid for non-integral
  * types.
  *
- * Also, the `key_type` and `value_type` must be the same size otherwise
- * the packed value will contain padding bits and accessing the packed
- * value would be undefined.
+ * Also, the `pair_type` must not contain any padding bits otherwise
+ * accessing the packed value would be undefined.
  *
  * @tparam pair_type The pair type that will be packed
  * @return true If the pair type can be packed