Fix concurrent_unordered_map to prevent accessing padding bits in pair_type

[davidwendt]

The concurrent_unordered_map class uses a thrust::pair type internally to manage key/value pairs. There is also an also an pair_packer utility that will pack the pair into a single type (uint32 or uint64) to better utilize the atomicCAS for storing the pair. The code also includes a path for cases when the thrust::pair is not packable.

The is_packable tests to make sure both pairs are integer types and can be packed in either a single uint32 or uint64 type. Unfortunately, if the types are different sizes (e.g. thrust::pair<uint64_t, uint32_t>) the pair-type also contains padding bits. This resulted in a crash in semi_left_join which resulted in a pair type of thrust::pair<uint32_t,bool>. Referencing the single packed value therefore included padding bits which is undefined behavior.

Debug session showing packed values with non-zero padding bits

[New Thread 0x7f05725fd700 (LWP 6311)]
[New Thread 0x7f0571c77700 (LWP 6312)]
[Switching focus to CUDA kernel 0, grid 5, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]

Thread 1 "map_test" hit Breakpoint 1, concurrent_unordered_map<int, bool, cudf::row_hasher<default_hash, true>, cudf::row_equality_comparator, default_allocator<thrust::pair<int, bool> > >::attempt_insert<thrust::pair<int, bool> > (this=0x7f0574fffb08, insert_location=0x7f0542e00448, insert_pair=0x7f0574fff9f8)
at /cudf/cpp/build/../src/hash/concurrent_unordered_map.cuh:283
283 if (old.packed == unused.packed) { return insert_result::SUCCESS; }
///// expected these 2 values to be equal but the random padding bits makes this false
(cuda-gdb) p old.packed
$1 = 6442450943
(cuda-gdb) p unused.packed
$2 = 8430733011321880575

(cuda-gdb) p old.pair.first
$3 = 2147483647
(cuda-gdb) p old.pair.second
$4 = true

(cuda-gdb) p unused.pair.first
$5 = 2147483647
(cuda-gdb) p unused.pair.second
$6 = true

(cuda-gdb) p/x unused.packed
$9 = 0x74fffb017fffffff
(cuda-gdb) p/x old.packed
$10 = 0x17fffffff

(cuda-gdb) p unused
$11 = {packed = 8430733011321880575, pair = {first = 2147483647, second = true}}
(cuda-gdb) p/x unused.pair.first
$12 = 0x7fffffff
(cuda-gdb) p/x unused.pair.second
$13 = 0x1
(cuda-gdb)

(cuda-gdb) p old
$14 = {packed = 6442450943, pair = {first = 2147483647, second = true}}
(cuda-gdb) p/x old.pair.first
$15 = 0x7fffffff
(cuda-gdb) p/x old.pair.second
$16 = 0x1
(cuda-gdb)

This PR solves this by adding a size check to the current is_packable(). The C++17 std::has_unique_object_representations is used here to ensure no padding bits are in the pair_type.

[vuule]

rerun tests

[codecov]

Codecov Report

❗ No coverage uploaded for pull request base (branch-21.08@606d854). Click here to learn what that means.
The diff coverage is n/a.

@@               Coverage Diff               @@
##             branch-21.08    #8348   +/-   ##
===============================================
  Coverage                ?   82.83%           
===============================================
  Files                   ?      109           
  Lines                   ?    17901           
  Branches                ?        0           
===============================================
  Hits                    ?    14828           
  Misses                  ?     3073           
  Partials                ?        0           

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 606d854...34851e9. Read the comment docs.

[harrism]

@gpucibot merge