Regarding erased_key_sentinel in static map

[sleeepyjack]

Opened on behalf of @avithemad.
Moving discussion from rapidsai/cudf#16734.

[sleeepyjack]

From @avithemad:

I understand the hash table better now based on your comments. But how I tested is as follows. I just inserted about 10k unique key value pairs (1-10000, 1-10000) identity key value in static map with 20k size. And did not delete anything so I assume for negative lookups it takes just one access.
I probe 60 million random values out of which only 3% values exist in hash table. What I observed was the kernel was taking about 70 milliseconds. But I implemented a very simple bloom filter to exit threads early if no match is found, and then kernel was taking about 15 ms. This is why I was doubting the negative lookups taking time.

The Bloom filter can probably fit into the L2$ so the lookup is considerably faster compared to the map lookup. It is a common practice to pre-filter the keys before performing the hash map lookup in cases where the matching rate is very low. So you did the right thing. BTW we do have a cuco::bloom_filter in the making (see #573) which can be directly used for this task. I'm currently finishing up this PR so it should be upstream in the near future. If you want, you can also checkout or merge the PR branch into you working branch and try it out yourself.

[avithemad]

I was suspecting the same thing, therefore I increased the size of the bloom filter. I initially used boolean mask, so runtime was optimal about 15ms, but just to test if bloom filter does not fit the L2 cache scenario I made each element in bloom filter 8 bytes as follows and runtime only increased to 20ms.
This is a very simple bloom filter implemented like this

long* bloomfilter
cudaMalloc(&bloomfilter, S_LEN * sizeof(long));

and S_LEN is about 800K, so total size of bloom filter is about 8MB.
Ideally hash table also is of similar size which I initialize like this

 auto supplier_map = cuco::static_map{S_LEN, cuco::empty_key{-1}, cuco::empty_value{-1}, thrust::equal_to<int>{},
                              cuco::double_hashing<TILE_SIZE, cuco::default_hash_function<int>>{}};

Would it may be because I use a simple identity hash function and static map's hash function is more general so takes more time to compute, but 70ms compared to 15ms is quite high right, computing the hash should be cheaper for integers? (Please pardon my ignorance)

Here are the 2 kernels I compared with and without bloom filter

template <typename Map>
__global__ void probe(Map part_map,
                      Map date_map,
                      Map supp_map,
                      int* lo_partkey,
                      int* lo_orderdate,
                      int* lo_suppkey,
                      int* lo_revenue,
                      int* d_year,
                      int* p_brand,
                      int* res,
                      int lo_size, long* suppkey_bloom, long* partkey_bloom)
{
  int tid = ((threadIdx.x + blockIdx.x * blockDim.x) / TILE_SIZE);
  if (tid >= lo_size) return;
  auto part_idx = part_map.find(lo_partkey[tid]);
  if (part_idx == part_map.end()) return;
  auto supp_idx = supp_map.find(lo_suppkey[tid]);
  if (supp_idx == supp_map.end()) return;
  auto date_idx = date_map.find(lo_orderdate[tid]);

and second with bloom

template <typename Map>
__global__ void probe(Map part_map,
                      Map date_map,
                      Map supp_map,
                      int* lo_partkey,
                      int* lo_orderdate,
                      int* lo_suppkey,
                      int* lo_revenue,
                      int* d_year,
                      int* p_brand,
                      int* res,
                      int lo_size, long* suppkey_bloom, long* partkey_bloom)
{
  int tid = ((threadIdx.x + blockIdx.x * blockDim.x) / TILE_SIZE);
  if (tid >= lo_size) return;
  if (suppkey_bloom[(lo_suppkey[tid]-1)%S_LEN] == false || partkey_bloom[(lo_partkey[tid]-1)%P_LEN] == false) return;
  auto part_idx = part_map.find(lo_partkey[tid]);
  if (part_idx == part_map.end()) return;
  auto supp_idx = supp_map.find(lo_suppkey[tid]);
  if (supp_idx == supp_map.end()) return;
  auto date_idx = date_map.find(lo_orderdate[tid]);

[sleeepyjack]

I just ran a benchmark on my L40s testing the lookup performance of static_multimap with varying key matching rate:

Benchmark Results

static_multimap_query_uniform_matching_rate

[0] NVIDIA L40S

Key	Value	Distribution	NumInputs	Occupancy	Multiplicity	MatchingRate	Samples	CPU Time	Noise	GPU Time	Noise	Elem/s
I32	I32	UNIFORM	100000000	0.5	1	0.1	13x	38.917 ms	0.04%	38.910 ms	0.04%	2.570G
I32	I32	UNIFORM	100000000	0.5	1	0.2	13x	39.483 ms	0.04%	39.477 ms	0.03%	2.533G
I32	I32	UNIFORM	100000000	0.5	1	0.3	13x	40.292 ms	0.02%	40.286 ms	0.02%	2.482G
I32	I32	UNIFORM	100000000	0.5	1	0.4	12x	42.685 ms	0.03%	42.678 ms	0.02%	2.343G
I32	I32	UNIFORM	100000000	0.5	1	0.5	12x	44.162 ms	0.03%	44.155 ms	0.02%	2.265G
I32	I32	UNIFORM	100000000	0.5	1	0.6	12x	44.987 ms	0.02%	44.980 ms	0.02%	2.223G
I32	I32	UNIFORM	100000000	0.5	1	0.7	12x	45.414 ms	0.03%	45.408 ms	0.02%	2.202G
I32	I32	UNIFORM	100000000	0.5	1	0.8	11x	45.618 ms	0.02%	45.612 ms	0.01%	2.192G
I32	I32	UNIFORM	100000000	0.5	1	0.9	11x	45.712 ms	0.04%	45.705 ms	0.04%	2.188G
I32	I32	UNIFORM	100000000	0.5	1	1	11x	45.777 ms	0.02%	45.771 ms	0.02%	2.185G
I64	I64	UNIFORM	100000000	0.5	1	0.1	11x	55.212 ms	0.02%	55.205 ms	0.02%	1.811G
I64	I64	UNIFORM	100000000	0.5	1	0.2	11x	56.879 ms	0.02%	56.873 ms	0.01%	1.758G
I64	I64	UNIFORM	100000000	0.5	1	0.3	11x	58.308 ms	0.02%	58.302 ms	0.02%	1.715G
I64	I64	UNIFORM	100000000	0.5	1	0.4	11x	66.309 ms	0.02%	66.302 ms	0.02%	1.508G
I64	I64	UNIFORM	100000000	0.5	1	0.5	11x	74.938 ms	0.01%	74.931 ms	0.01%	1.335G
I64	I64	UNIFORM	100000000	0.5	1	0.6	11x	83.604 ms	0.02%	83.597 ms	0.01%	1.196G
I64	I64	UNIFORM	100000000	0.5	1	0.7	11x	92.272 ms	0.01%	92.265 ms	0.01%	1.084G
I64	I64	UNIFORM	100000000	0.5	1	0.8	11x	100.935 ms	0.01%	100.928 ms	0.01%	990.809M
I64	I64	UNIFORM	100000000	0.5	1	0.9	11x	109.575 ms	0.01%	109.568 ms	0.01%	912.674M
I64	I64	UNIFORM	100000000	0.5	1	1	11x	118.246 ms	0.02%	118.239 ms	0.02%	845.747M

As you can see, the lookup performance is even higher if the matching rate is low.
You can reproduce these numbers yourself using the following steps:

cd cuCollections
ci/build.sh -b
build/latest/benchmarks/STATIC_MULTIMAP_BENCH -b 11

If you want to play around with the benchmark axes for e.g. Occupancy or NumInputs then this is also possible. Consult build/latest/benchmarks/STATIC_MULTIMAP_BENCH -h which provides some examples on how to pass parameter values to the benchmark setup.

Would it may be because I use a simple identity hash function

Depending on your input key distribution, this could certainly impact performance. Let's say your keys are a dense contiguous sequence (sorted or not), then with identity_hash all the keys would end up clumped in the same region of the map, hindering the early exit from being triggered as there are very few empty slots in this region. What you want is a uniform slot occupancy across the map, so the chance of finding an empty slot early is as high as possible. From my experience, the computational load of computing a (proper) hash function doesn't have a significant impact on the end-to-end performance. You could give cuco::xxhash_32 a spin and see how it performs.

Another factor that might impact performance is the load factor of the map. The more occupied the map is, the lower the chance of finding an empty slot. What's the map's occupancy for your current setup?

[avithemad]

You are right! I just created an empty hash table and probed it with 60M rows, it was around 50ms, which is 20 ms lesser than if it was occupied.

My setup is :

• load factor: 2, hash table is twice the size of keys. So I initialize the capacity of static_map with 2*800,000
• occupancy: 0.03, about 32,882/800,000 key, value actually exist in hash map
• All keys are contiguous from 1-800,000 (and we only inserted selected tuples in the map).

But even with the empty hash table it looks like it is taking lot of time compared to simple implementations.

My main goal was to implement the SSB queries(specifically q21) that were written with CRYSTAL library using the cuco library.
https://github.com/anilshanbhag/crystal/blob/master/src/ssb/q21.cu
So I am using the same inputs that this code uses, but just replace hash map implementation with cuco. Their simple implementation of hash table makes probe phase finish in 14ms (they assume perfect hashing, with identity hash function), so I was expecting the empty probe lookup also to take similar time but it takes 50ms.
q21.txt
I have attached the same implementation with cuco library, q21.txt. But with line 59 uncommented similar time to crystal is obtained, which I think should be achievable with cuco::bloom_filter. I will try that out.

By the way here is the link to the paper for details: https://dl.acm.org/doi/pdf/10.1145/3318464.3380595

[avithemad]

I think I got the reason. The hash table they implement in crystal is just 1D array similar to a bloom filter, therefore they get comparable performance. I think this performance of empty probe is reasonable assuming uniform distribution of probing.
Thank you for clarifying about the performance.

[sleeepyjack]

(they assume perfect hashing, with identity hash function)

Oh, I see now. Perfect hashing in combination with an identity hasher basically boils down to map[key] = value aka there is no probing involved whatsoever since the mapping is bijective by-design. The reason you're seeing the performance degradation is because cuco cannot introspect the data to check if probing is required or not. So the probing logic is always in the kernel which gives you worse codegen compared to the simple lookup table/ Bloom filter approach. So the regression is expected.

We tried to add perfect hashing as a special case to our implementation but it's still a work in progress. One could argue that with perfect hashing + identity hashing you don't even need any of the key benefits a hash map provides which is the ability to map a sparsely populated but huge key domain (think of uint64_t which has 2^64 possible values) onto a dense memory domain. There's no way you could allocate a lookup table providing slots for all possible uint64_t on any architecture in the world (even combined). Hash map probing implements an injective mapping that solves this problem. In your case, the input domain is already dense and fits nicely into DRAM. Thus, it's a different data structure with a disjunct application field.

I'd say you go with the simple Bloom filter approach for best performance.

Let me know if you have more questions; otherwise, feel free to mark this discussion as resolved.

Thanks!

[avithemad]

@sleeepyjack , I think I cannot resolve this discussion (Don't see any option to do so).
I do not have anymore question on this for now, I think if you can please resolve the discussion.

Thank you.

[sleeepyjack]

Whups, you're right...I totally forgot that I opened this discussion xD