Count method in static_multimap returning 0

[sarda-devesh]

Hello,

I was trying to utilize a static multimap and get the count for various keys using device references to the static multimap. I was consistently getting that the count method was returning 0 regardless of the key that I was using. Interestingly, I am also consistently getting that the contains was returning true when searching for those same keys. Thus, I wrote the following test program:

template <uint32_t cg_size, typename InsertView>
__global__ void map_filler(InsertView insert_view, int num_items) {
    auto tile = cooperative_groups::tiled_partition<cg_size>(cooperative_groups::this_thread_block());
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while(tid < num_items) {
      insert_view.insert(tile, {tid, tid});
      tid += gridDim.x * blockDim.x;
    }
}

template <uint32_t cg_size, typename CountView>
__global__ void count_checker(CountView count_view, int num_items) {
    auto tile = cooperative_groups::tiled_partition<cg_size>(cooperative_groups::this_thread_block());
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while(tid < num_items) {
      if(count_view.contains(tile, tid) && count_view.count(tile, tid) == 0) {
        printf("Tid %d exists in map but has count 0\n", tid);
      }
      tid += gridDim.x * blockDim.x;
    }
}

int main(int argc, char** argv) {
    high_resolution_clock::time_point start;
    high_resolution_clock::time_point end;
    duration<double, std::milli> duration_ms;

    int num_items = 6001215; int threads_per_block = 1024;
    std::size_t capacity = 2 * num_items;
    cuco::static_multimap<int, int> test_map{capacity, cuco::empty_key{-1}, cuco::empty_value{-1}};
    auto constexpr cg_size = test_map.cg_size();

    // Insert into the map
    start = high_resolution_clock::now();
    int num_blocks = (num_items + threads_per_block - 1)/(10 * threads_per_block);
    map_filler<cg_size><<<num_blocks, threads_per_block>>>(test_map.get_device_mutable_view(), num_items);
    cudaDeviceSynchronize();
    end = high_resolution_clock::now();
    duration_ms = std::chrono::duration_cast<duration<double, std::milli>>(end - start);
    std::cout << "Insert time of " << duration_ms.count() << " ms" << std::endl;

    // Query the map
    start = high_resolution_clock::now();
    count_checker<cg_size><<<num_blocks, threads_per_block>>>(test_map.get_device_view(), num_items);
    cudaDeviceSynchronize();
    end = high_resolution_clock::now();
    duration_ms = std::chrono::duration_cast<duration<double, std::milli>>(end - start);
    std::cout << "Count time of " << duration_ms.count() << " ms" << std::endl;
}

Running this programs outputs:

Tid 75707 exists in map but has count 0
Tid 75708 exists in map but has count 0
Tid 75709 exists in map but has count 0
Tid 75710 exists in map but has count 0
Tid 75711 exists in map but has count 0
Tid 57089 exists in map but has count 0
Tid 57090 exists in map but has count 0
Tid 57091 exists in map but has count 0
Tid 57092 exists in map but has count 0
Tid 57093 exists in map but has count 0

This indicates that when using the reference returned by get_device_view, the keys seem to be present in the map when using contains but don't seem to present in the map when using count

[PointKernel]

@sarda-devesh Thanks for reporting the issue.

I didn't find anything suspicious with a glance at your sample code. Let me test it locally and get back to you later.

[sarda-devesh]

As a follow up, I updated the code as such:

template <uint32_t cg_size, typename InsertView>
__global__ void map_filler(InsertView insert_view, int num_items) {
    auto tile = cooperative_groups::tiled_partition<cg_size>(cooperative_groups::this_thread_block());
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while(tid < num_items) {
      printf("Inserting for tid %d\n", tid);  
      insert_view.insert(tile, {tid, tid});
      tid += gridDim.x * blockDim.x;
    }
}

template <uint32_t cg_size, typename CountView>
__global__ void count_checker(CountView count_view, int num_items) {
    auto tile = cooperative_groups::tiled_partition<cg_size>(cooperative_groups::this_thread_block());
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while(tid < num_items) {
      if(count_view.contains(tile, tid) && count_view.count(tile, tid) == 0) {
        printf("Tid %d exists in map but has count 0\n", tid);
      }

      tid += gridDim.x * blockDim.x;
    }
}

int main(int argc, char** argv) {
    high_resolution_clock::time_point start;
    high_resolution_clock::time_point end;
    duration<double, std::milli> duration_ms;

    int num_items = 2; int threads_per_block = 1024; 
    std::size_t capacity = 2 * num_items;
    cuco::static_multimap<int, int> test_map{capacity, cuco::empty_key{-1}, cuco::empty_value{-1}};
    auto constexpr cg_size = test_map.cg_size();

    // Insert into the map
    start = high_resolution_clock::now();
    int num_blocks = (num_items + threads_per_block - 1)/threads_per_block;
    map_filler<cg_size><<<num_blocks, threads_per_block>>>(test_map.get_device_mutable_view(), num_items);
    cudaDeviceSynchronize();
    end = high_resolution_clock::now();
    duration_ms = std::chrono::duration_cast<duration<double, std::milli>>(end - start);
    std::cout << "Insert time of " << duration_ms.count() << " ms" << std::endl;

    // Query the map
    start = high_resolution_clock::now();
    count_checker<cg_size><<<num_blocks, threads_per_block>>>(test_map.get_device_view(), num_items);
    cudaDeviceSynchronize();
    end = high_resolution_clock::now();
    duration_ms = std::chrono::duration_cast<duration<double, std::milli>>(end - start);
    std::cout << "Count time of " << duration_ms.count() << " ms" << std::endl;
}

which produces the output:

Inserting for tid 0
Inserting for tid 1
Insert time of 2.25478 ms
Tid 1 exists in map but has count 0
Count time of 0.192259 ms

Thus, it looks like count returns a non-zero value for tid 0 but returns 0 for tid 1. I am not sure what the cause is but some interesting behaviour that I noticed.

[PointKernel]

OK, so there are two issues in the code:

1. As mentioned by @sleeepyjack, the CG count function returns a per-thread count thus it's users' responsibility to sum up all threads within a CG to get the final per-CG count.
2. The idea of CG insert/query is to utilize multiple threads performing hash map operations for one input key element. In your sample code, e.g., when map cg_size == 8 (default value for multimap), we are supposed to use the first tile (thread 0, 1, 2, ... 7) to insert the first pair {0, 0} and the second tile (thread 8, 9, 10, ... 15) to insert {1, 1} as opposed to thread 0 of tile 0 inserting {0, 0} and thread 1 of tile 0 inserting {1, 1}. That's an undefined behavior.

As explained, a potential solution to your use case https://godbolt.org/z/38vMY9669

[sleeepyjack]

I think the behavior you're seeing here arises from mixin CG and per-thread computation. The docs of the CG count functions state "[returns] Number of matches found by the current thread" (here). So in order to get the number of total occurrences of a key, you have to aggregate the count across all threads in the CG first, e.g. by using tile.reduce(...) (see docs).

[sarda-devesh]

First of all, thank you so much for your quick and helpful response. @PointKernel The example that you linked was really helpful. I just want to confirm my understanding of something: If I define my static_multimap as such:

using probe = cuco::legacy::linear_probing<1, cuco::default_hash_function<int>>;
  cuco::static_multimap<int, int, cuda::thread_scope_device, cuco::cuda_allocator<char>, probe> test_map{capacity, cuco::empty_key{-1}, cuco::empty_value{-1}};

then the cg_size will be 1 which I have verified through printing it out. In that case tile_id == tid and thus I will be able to get the desired behavior of each thread being responsible for inserting one pair and calling the count function will give me the total count for that key.

Once again, thank you so much for your help.

[PointKernel]

In that case tile_id == tid and thus I will be able to get the desired behavior of each thread being responsible for inserting one pair and calling the count function will give me the total count for that key.

That's correct.