|
17 | 17 | #include <cuco/static_map.cuh> |
18 | 18 |
|
19 | 19 | #include <cuda/std/array> |
| 20 | +#include <cuda/std/limits> |
20 | 21 |
|
21 | 22 | #include <cooperative_groups.h> |
22 | 23 |
|
|
36 | 37 | */ |
37 | 38 |
|
38 | 39 | // Basic types |
39 | | -using Key = int; |
40 | 40 | using Value = uint32_t; |
| 41 | +using Key = int; |
41 | 42 |
|
42 | 43 | // Sentinel values for empty slots |
43 | 44 | Key constexpr empty_key_sentinel = -1; |
44 | | -Value constexpr empty_value_sentinel = 0xffffffff; |
| 45 | +Value constexpr empty_value_sentinel = cuda::std::numeric_limits<Value>::min(); |
45 | 46 |
|
46 | 47 | // Map configuration |
47 | 48 | std::size_t constexpr capacity = 100'000; |
@@ -126,18 +127,18 @@ int main() |
126 | 127 | { |
127 | 128 | // Step 1: Initialize the pre-allocated storage |
128 | 129 | init_kernel<<<1, 128>>>(); |
129 | | - cudaDeviceSynchronize(); |
| 130 | + CUCO_CUDA_TRY(cudaDeviceSynchronize()); |
130 | 131 |
|
131 | 132 | // Step 2: Insert some key-value pairs |
132 | 133 | insert_kernel<<<2, 32>>>(); |
133 | | - cudaDeviceSynchronize(); |
| 134 | + CUCO_CUDA_TRY(cudaDeviceSynchronize()); |
134 | 135 |
|
135 | 136 | // Step 3: Find and verify the inserted pairs |
136 | 137 | find_kernel<<<2, 32>>>(); |
137 | 138 |
|
138 | 139 | // Check results - expect to find all 64 keys (2 blocks * 32 threads) |
139 | 140 | int host_found_count; |
140 | | - cudaMemcpyFromSymbol(&host_found_count, found_count, sizeof(int)); |
| 141 | + CUCO_CUDA_TRY(cudaMemcpyFromSymbol(&host_found_count, found_count, sizeof(int))); |
141 | 142 | int expected_count = 2 * 32; // Total number of keys inserted and queried |
142 | 143 |
|
143 | 144 | if (host_found_count == expected_count) { |
|
0 commit comments