The static_map of custom type becomes very slow after modifying the custom_key_equal function.

[Po1sson]

As a key with two member variables, it is reasonable to require both values ​​to match when matching the key. In large-scale hash tables, it is common for many keys to have the same first member variable but different second member variables. However, I modified custom_key_equal to require that both member variables are equal means that the key is equal, and the insertion speed of static_map became so slow that I could not wait until the end of the program execution.
Is this normal behavior? If so, how can I optimize it?

I modified the key generation rules so that the first member variable key.a of some keys is the same, but key.b is different to distinguish them.

auto pairs_begin = thrust::make_transform_iterator(
thrust::make_counting_iterator<int32_t>(0),
cuda::proclaim_return_type<cuco::pair<custom_key_type, custom_value_type>>(
[] __device__(auto i) {
return cuco::pair{ custom_key_type{i%100,i}, custom_value_type{i} };
}));

The modified function for judging key equality is as follows.

struct custom_key_equal {
    __device__ bool operator()(custom_key_type const& lhs, custom_key_type const& rhs) const noexcept
    {
        return ((lhs.a == rhs.a) && (lhs.b == rhs.b));
    }
};

The complete code is as follows:

#include <cuco/static_map.cuh>

#include <cuda/functional>
#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/logical.h>
#include <thrust/transform.h>

 // User-defined key type
struct custom_key_type {
    int32_t a;
    int32_t b;

    __host__ __device__ custom_key_type() {}
    __host__ __device__ custom_key_type(int32_t x) : a{ x }, b{ x } {}
    __host__ __device__ custom_key_type(int32_t x, int32_t y) : a{ x }, b{ y } {}
};

// User-defined value type
struct custom_value_type {
    int32_t f;
    int32_t s;

    __host__ __device__ custom_value_type() {}
    __host__ __device__ custom_value_type(int32_t x) : f{ x }, s{ x } {}
};

// User-defined device hash callable
struct custom_hash {
    __device__ uint32_t operator()(custom_key_type const& k) const noexcept { return k.a; };
};

// User-defined device key equal callable
struct custom_key_equal {
    __device__ bool operator()(custom_key_type const& lhs, custom_key_type const& rhs) const noexcept
    {
        return ((lhs.a == rhs.a) && (lhs.b == rhs.b));
    }
};

int main(void)
{
    constexpr std::size_t num_pairs = 5000'000;

    // Set emtpy sentinels
    auto const empty_key_sentinel = custom_key_type{ -1 };
    auto const empty_value_sentinel = custom_value_type{ -1 };

    // Create an iterator of input key/value pairs
    auto pairs_begin = thrust::make_transform_iterator(
        thrust::make_counting_iterator<int32_t>(0),
        cuda::proclaim_return_type<cuco::pair<custom_key_type, custom_value_type>>(
            [] __device__(auto i) {
        return cuco::pair{ custom_key_type{i%100,i}, custom_value_type{i} };
    }));

    // Construct a map with 100,000 slots using the given empty key/value sentinels. Note the
    // capacity is chosen knowing we will insert 80,000 keys, for an load factor of 80%.
    auto map = cuco::static_map{ cuco::extent<std::size_t, 6500'000>{},
                                cuco::empty_key{empty_key_sentinel},
                                cuco::empty_value{empty_value_sentinel},
                                custom_key_equal{},
                                cuco::linear_probing<1, custom_hash>{} };

    // Inserts 80,000 pairs into the map by using the custom hasher and custom equality callable
    map.insert(pairs_begin, pairs_begin + num_pairs);

    // Reproduce inserted keys
    auto insert_keys =
        thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
            cuda::proclaim_return_type<custom_key_type>(
                [] __device__(auto i) { return custom_key_type{ i % 100,i }; }));

    thrust::device_vector<bool> contained(num_pairs);

    // Determine if all the inserted keys can be found by using the same hasher and equality
    // function as `insert`. If a key `insert_keys[i]` doesn't exist, `contained[i] == false`.
    map.contains(insert_keys, insert_keys + num_pairs, contained.begin());
    // This will fail due to inconsistent hash and key equal.
    // map.contains(insert_keys, insert_keys + num_pairs, contained.begin());

    // All inserted keys are contained
    auto const all_contained =
        thrust::all_of(contained.begin(), contained.end(), [] __device__(auto const& b) { return b; });
    if (all_contained) { std::cout << "Success! Found all values.\n"; }

    return 0;
}```

[PointKernel]

auto pairs_begin = thrust::make_transform_iterator(
thrust::make_counting_iterator<int32_t>(0),
cuda::proclaim_return_type<cuco::pair<custom_key_type, custom_value_type>>(
[] __device__(auto i) {
return cuco::pair{ custom_key_type{i%100,i}, custom_value_type{i} };
}));

By doing so, you are adding more duplicates to the input. Thus more hash collisions. To solve the problem, you may want to update your hash function accordingly to minimize collisions, e.g.

// User-defined device hash callable
struct custom_hash {
    __device__ uint32_t operator()(custom_key_type const& k) const noexcept { return k.a + k.b; };
};

Or hash your struct properly like https://stackoverflow.com/questions/19195183/how-to-properly-hash-the-custom-struct

[Po1sson]

Thank you for your response. After comparing with the solution you provided, I reviewed my code and found that there were some deviations in the code I uploaded yesterday. However, your solution helped me identify the issue.
In fact, I imitated the example and designed a kernel function for inserting data into static_map.
My data is stored in a continuous memory, the pointer is data, and my expectation is that each thread is responsible for reading four consecutive values, the first two values ​​are used as keys, and the last two values ​​are used as values.
But I found that this would be very slow. It takes more than 10 seconds to insert 1,000,000 key-value pairs on 2080Ti.
map_ref.insert(cuco::pair{ point(data[tid * 4],data[tid * 4 + 1]), point(data[tid * 4 + 2],data[tid * 4 + 3]) });
But if I use tid to replace a value in the key, the time consumption will be reduced to about 1ms. In addition, if I use tid to replace the first member variable of the key and replace the second member variable, the insertion time is different.

It takes about 1.5ms to call the following statement to insert 1000000 key-value pairs：
map_ref.insert(cuco::pair{ point(data[tid * 4],tid), point(data[tid * 4 + 2],data[tid * 4 + 3]) });

It takes about 1.16ms to call the following statement to insert 1000000 key-value pairs：
map_ref.insert(cuco::pair{ point(data[tid * 4 + 1],tid), point(data[tid * 4 + 2],data[tid * 4 + 3]) });

It takes about 1.03ms to call the following statement to insert 1000000 key-value pairs：
map_ref.insert(cuco::pair{ point(tid,data[tid * 4]), point(data[tid * 4 + 2],data[tid * 4 + 3]) });

It takes about 10466.2ms to call the following statement to insert 1000000 key-value pairs：
map_ref.insert(cuco::pair{ point(data[tid * 4],data[tid * 4 + 1]), point(data[tid * 4 + 2],data[tid * 4 + 3]) });

Even if the following method is used to insert key-value pairs, it takes about 8600ms.

    auto pairs_begin = thrust::make_transform_iterator(
        D.begin(),
        cuda::proclaim_return_type<cuco::pair<point, point>>(
            [] __device__(auto i) {
        return cuco::pair{ point{i.a,i.b}, point{i.c,i.d} };
    }));

    map.insert(pairs_begin, pairs_begin + num_pairs);

The complete code is as follows，I have no idea where the problem is, I would really appreciate it if you could help me figure it out.

#include <iostream>
#include <cuda_runtime.h>
#include <fstream>
#include <cuco/static_map.cuh>
#include <cuda/functional>
#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/logical.h>
#include <thrust/transform.h>
#include <thrust/host_vector.h>
#include <vector>
#include <chrono>

struct segment {
    int32_t a;
    int32_t b;
    int32_t c;
    int32_t d;
    __host__ __device__ segment() {}
    __host__ __device__ segment(int32_t a, int32_t b, int32_t c, int32_t d) : a{ a }, b{ b }, c{ c }, d{ d } {}
};
struct point {
    int32_t a;
    int32_t b;

    __host__ __device__ point() {}
    __host__ __device__ point(int32_t x) : a{ x }, b{ x } {}
    __host__ __device__ point(int32_t x, int32_t y) : a{ x }, b{ y } {}
};

struct custom_hash {
    __device__ uint32_t operator()(point const& k) const noexcept {
        int32_t hash = 17;
        hash = hash * 31 + static_cast<int32_t>(k.a);
        hash = hash * 31 + static_cast<int32_t>(k.b);
        return hash;
    };
};

struct custom_key_equal {
    __device__ bool operator()(point const& lhs, point const& rhs) const noexcept
    {
        return ((lhs.a == rhs.a) && (lhs.b == rhs.b));
    }
};

template <typename Map, typename Datatype>
__global__ void filtered_insert(Map map_ref, std::size_t num_keys, Datatype* data, int* num_inserted)
{
    auto tid = threadIdx.x + blockIdx.x * blockDim.x;
    std::size_t counter = 0;
    while (tid < num_keys) {
        //if (map_ref.insert(cuco::pair{ point(data[tid * 4],data[tid * 4 + 1]), point(data[tid * 4 + 2],data[tid * 4 + 3]) })) {//10466.2
        if (map_ref.insert(cuco::pair{ point(data[tid * 4],tid), point(data[tid * 4 + 2],data[tid * 4 + 3]) })) { //1.53798
        //if (map_ref.insert(cuco::pair{ point(tid,data[tid * 4 + 1]), point(data[tid * 4 + 2],data[tid * 4 + 3]) })) { //1.02288
            counter++;
        }
        tid += gridDim.x * blockDim.x;
    }
    atomicAdd(num_inserted, counter);
}

thrust::host_vector<segment> readData() {
    thrust::host_vector<segment> data;
    std::ifstream inFile("data6923200.txt");
    if (!inFile) {
        std::cerr << "Unable to open file for reading" << std::endl;
        return data;
    }
    int a, b, c, d, e;
    while (inFile >> a >> b >> c >> d >> e) {
        int dir = a == c ? -1 * e : e;
        if (dir == 1) {
            data.push_back(segment(a, b, c, d));
        }
        else {
            data.push_back(segment(c, d, a, b));
        }
    }
    inFile.close();
    return data;
}

int main(void)
{
    auto H = readData();
    std::size_t const num_pairs = 1000000;
    std::cout << "Num Pairs:" << num_pairs << std::endl;

    auto const empty_key_sentinel = point{ -1 };
    auto const empty_value_sentinel = point{ -1 };

    thrust::device_vector<segment> D = H;

    std::size_t const capacity = std::ceil(num_pairs / 0.7);
    auto map = cuco::static_map{ capacity,
                                cuco::empty_key{empty_key_sentinel},
                                cuco::empty_value{empty_value_sentinel},
                                custom_key_equal{},
                                cuco::linear_probing<1, custom_hash>{} };

    auto insert_ref = map.ref(cuco::insert);
    thrust::device_vector<int> num_inserted(1);
    thrust::device_vector<int> pool_flag(1);

    auto constexpr block_size = 512;
    auto const grid_size = 128;

    segment* raw_ptr = thrust::raw_pointer_cast(D.data());
    int* address_of_a = &(raw_ptr[0].a);

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start);

    filtered_insert << <grid_size, block_size >> > (insert_ref,
        (std::size_t)1000000,
        address_of_a,
        num_inserted.data().get());

    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    std::cout << "Kernel execution time: " << milliseconds << " ms" << std::endl;

    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    std::cout << "Number of keys inserted: " << num_inserted[0] << std::endl;
    return 0;
}

[Po1sson]

I re-tested and replaced the second number of each data key-value pair, that is, the second member of the key, with a number that grows from zero. I found that after this modification, it only took less than 2ms to insert 1,000,000 key-value pairs. Therefore, I judged that the performance loss was not caused by accessing global memory, but because the hash function was not suitable, resulting in a large number of hash collisions.
Thanks again for your help.

[sleeepyjack]

the insertion speed of static_map became so slow that I could not wait until the end of the program execution

Did the operation finish at all or are we running into an infinite (probing) loop? Briefly reading through your code, I can't spot anything wrong with it apart from what @PointKernel already mentioned.

We had some recent changes/bugfixes to our insertion backend. Can you pull the latest dev changes and see if it fixes the problem?

[Po1sson]

Thank you for your response. Yesterday, I made a silly mistake while uploading the code. Your answer made me realize that the execution efficiency of my code was hindered by access to global memory(I guess). My data reaches hundreds of megabytes, so it is stored in global memory. I am currently exploring and trying to solve this issue.

[Po1sson]

I seem to have found the problem.
I re-tested and replaced the second number of each data key-value pair, that is, the second member of the key, with a number that increases from zero. I found that after this modification, it only takes less than 2ms to insert 1,000,000 key-value pairs. Therefore, I judge that the performance loss is not due to the access to global memory, but because the hash function is not suitable, resulting in a large number of hash collisions.
Thanks again for your help.