diff --git a/examples/cpp-console-cuco-static-map-bench/CMakeLists.txt b/examples/cpp-console-cuco-static-map-bench/CMakeLists.txt new file mode 100644 index 00000000..ee15e9ad --- /dev/null +++ b/examples/cpp-console-cuco-static-map-bench/CMakeLists.txt @@ -0,0 +1,56 @@ +cmake_minimum_required(VERSION 3.20) + +project(cpp-console-cuco-static-map-bench LANGUAGES CXX CUDA) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CUDA_STANDARD 17) +set(CMAKE_CUDA_STANDARD_REQUIRED ON) +set(CMAKE_CUDA_EXTENSIONS OFF) + +set( + CUCO_ROOT + "" + CACHE PATH + "Path to a local cuCollections checkout." +) + +find_package(CUDAToolkit REQUIRED) + +find_path( + CUCO_INCLUDE_DIR + NAMES cuco/static_map.cuh + HINTS ${CUCO_ROOT} + PATH_SUFFIXES include + REQUIRED +) + +add_executable(cpp-console-cuco-static-map-bench + src/main.cu +) + +target_include_directories( + cpp-console-cuco-static-map-bench + PRIVATE + "${CUCO_INCLUDE_DIR}" +) + +target_link_libraries( + cpp-console-cuco-static-map-bench + PRIVATE + CUDA::cudart +) + +target_compile_options( + cpp-console-cuco-static-map-bench + PRIVATE + $<$:--expt-extended-lambda> +) + +if(APPLE OR UNIX) + set_property( + TARGET cpp-console-cuco-static-map-bench + PROPERTY BUILD_RPATH "$" + ) +endif() diff --git a/examples/cpp-console-cuco-static-map-bench/README.md b/examples/cpp-console-cuco-static-map-bench/README.md new file mode 100644 index 00000000..964154d6 --- /dev/null +++ b/examples/cpp-console-cuco-static-map-bench/README.md @@ -0,0 +1,44 @@ +# cpp-console-cuco-static-map-bench + +Standalone `cuCollections` baseline for the same raw `uint64_t` `.keys` domains +used by the PerfectHash GPU experiments. + +This benchmark: + +1. loads a `.keys` file as `uint64_t` keys, +2. builds a `cuco::static_map` mapping `key -> ordinal`, +3. runs `find_async()` over the same key stream, +4. reports build and lookup timings, including normalized `ns/key`. + +This is the closest simple `cuCollections` baseline to the current PerfectHash +GPU key-to-index lookup path. + +## Build + +```bash +cmake -S examples/cpp-console-cuco-static-map-bench \ + -B build/examples/cpp-console-cuco-static-map-bench \ + -DCUCO_ROOT=/path/to/cucollections + +cmake --build build/examples/cpp-console-cuco-static-map-bench -j +``` + +## Run + +```bash +./build/examples/cpp-console-cuco-static-map-bench/cpp-console-cuco-static-map-bench \ + --keys-file /mnt/data/tpch/scale-10/perfecthash-keys/c_custkey_q03_ph-64.keys \ + --load-factor 0.5 +``` + +Optional arguments: + +- `--keys-file ` +- `--max-keys ` +- `--load-factor ` +- `--device ` +- `--warmup ` +- `--iterations ` +- `--csv` +- `--csv-header` +- `--no-verify` diff --git a/examples/cpp-console-cuco-static-map-bench/src/main.cu b/examples/cpp-console-cuco-static-map-bench/src/main.cu new file mode 100644 index 00000000..e97bd95a --- /dev/null +++ b/examples/cpp-console-cuco-static-map-bench/src/main.cu @@ -0,0 +1,471 @@ +#include +#include + +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +using key_type = std::uint64_t; +using value_type = std::uint32_t; +using pair_type = cuco::pair; + +struct options { + std::string keys_file; + std::string probe_keys_file; + std::uint64_t max_keys = 0; + std::uint64_t max_probe_keys = 0; + double load_factor = 0.5; + int device_ordinal = 0; + int warmup = 2; + int iterations = 10; + bool csv = false; + bool csv_header = false; + bool verify = true; +}; + +struct result_row { + std::string key_source; + std::string probe_source; + std::string device_name; + std::size_t build_key_count = 0; + std::size_t build_key_bytes = 0; + std::size_t probe_key_count = 0; + std::size_t probe_key_bytes = 0; + std::size_t pair_bytes = 0; + std::size_t output_bytes = 0; + std::size_t slot_bytes = 0; + std::size_t capacity = 0; + std::size_t size = 0; + double requested_load_factor = 0.0; + double actual_load_factor = 0.0; + double h2d_keys_ms = 0.0; + double pair_prepare_ms = 0.0; + double insert_ms = 0.0; + double build_ms = 0.0; + double lookup_avg_ms = 0.0; + double lookup_min_ms = 0.0; + double lookup_max_ms = 0.0; + double lookup_ns_per_key = 0.0; + double verify_ms = 0.0; + std::size_t vram_buffers_bytes = 0; + std::size_t vram_total_bytes = 0; + int device_ordinal = 0; +}; + +struct cuda_stream_handle { + cudaStream_t value{}; + ~cuda_stream_handle() + { + if (value != nullptr) { cudaStreamDestroy(value); } + } +}; + +struct cuda_event_handle { + cudaEvent_t value{}; + ~cuda_event_handle() + { + if (value != nullptr) { cudaEventDestroy(value); } + } +}; + +void throw_if_cuda(cudaError_t status, char const* what) +{ + if (status != cudaSuccess) { + throw std::runtime_error(std::string(what) + " failed: " + cudaGetErrorString(status)); + } +} + +double ms_between(cudaEvent_t start, cudaEvent_t stop) +{ + float ms = 0.0f; + throw_if_cuda(cudaEventElapsedTime(&ms, start, stop), "cudaEventElapsedTime"); + return static_cast(ms); +} + +double ns_per_key(double milliseconds, std::size_t key_count) +{ + if (key_count == 0) { return 0.0; } + return (milliseconds * 1'000'000.0) / static_cast(key_count); +} + +std::vector load_keys_file(std::string const& path, std::uint64_t max_keys) +{ + std::ifstream input(path, std::ios::binary | std::ios::ate); + if (!input) { throw std::runtime_error("Unable to open keys file: " + path); } + + auto const size = input.tellg(); + if (size < 0) { throw std::runtime_error("Unable to determine keys file size: " + path); } + if ((static_cast(size) % sizeof(key_type)) != 0) { + throw std::runtime_error("Keys file size is not a multiple of 8 bytes: " + path); + } + + auto const available_keys = + static_cast(size) / static_cast(sizeof(key_type)); + auto const keys_to_read = + (max_keys != 0 && max_keys < available_keys) ? max_keys : available_keys; + + std::vector keys(static_cast(keys_to_read)); + input.seekg(0, std::ios::beg); + input.read(reinterpret_cast(keys.data()), + static_cast(keys.size() * sizeof(key_type))); + if (!input) { throw std::runtime_error("Unable to read keys file contents: " + path); } + + return keys; +} + +void print_usage(char const* argv0) +{ + std::cout << "Usage: " << argv0 + << " --keys-file [--probe-keys-file ] [--max-keys ] [--max-probe-keys ] [--load-factor ]\n" + " [--device ] [--warmup ] [--iterations ]\n" + " [--csv] [--csv-header] [--no-verify]\n"; +} + +options parse_args(int argc, char** argv) +{ + options opts; + for (int index = 1; index < argc; ++index) { + std::string arg = argv[index]; + if (arg == "--help" || arg == "-h") { + print_usage(argv[0]); + std::exit(0); + } + if (arg == "--keys-file" && index + 1 < argc) { + opts.keys_file = argv[++index]; + continue; + } + if (arg == "--probe-keys-file" && index + 1 < argc) { + opts.probe_keys_file = argv[++index]; + continue; + } + if (arg == "--max-keys" && index + 1 < argc) { + opts.max_keys = std::stoull(argv[++index]); + continue; + } + if (arg == "--max-probe-keys" && index + 1 < argc) { + opts.max_probe_keys = std::stoull(argv[++index]); + continue; + } + if (arg == "--load-factor" && index + 1 < argc) { + opts.load_factor = std::stod(argv[++index]); + continue; + } + if (arg == "--device" && index + 1 < argc) { + opts.device_ordinal = std::stoi(argv[++index]); + continue; + } + if (arg == "--warmup" && index + 1 < argc) { + opts.warmup = std::stoi(argv[++index]); + continue; + } + if (arg == "--iterations" && index + 1 < argc) { + opts.iterations = std::stoi(argv[++index]); + continue; + } + if (arg == "--csv") { + opts.csv = true; + continue; + } + if (arg == "--csv-header") { + opts.csv = true; + opts.csv_header = true; + continue; + } + if (arg == "--no-verify") { + opts.verify = false; + continue; + } + throw std::runtime_error("Unknown argument: " + arg); + } + + if (opts.keys_file.empty()) { throw std::runtime_error("--keys-file is required"); } + if (!(opts.load_factor > 0.0 && opts.load_factor < 1.0)) { + throw std::runtime_error("--load-factor must be in the open interval (0, 1)"); + } + if (opts.warmup < 0) { throw std::runtime_error("--warmup must be >= 0"); } + if (opts.iterations <= 0) { throw std::runtime_error("--iterations must be > 0"); } + + return opts; +} + +__global__ void fill_pairs_kernel(key_type const* keys, pair_type* pairs, std::size_t count) +{ + auto const tid = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (tid < count) { pairs[tid] = pair_type{keys[tid], static_cast(tid)}; } +} + +void print_csv_header() +{ + std::cout << "key_source,build_key_count,build_key_bytes,probe_source,probe_key_count,probe_key_bytes,pair_bytes,output_bytes,slot_bytes,requested_load_factor," + "capacity,size,actual_load_factor,h2d_keys_ms,pair_prepare_ms,insert_ms,build_ms," + "lookup_avg_ms,lookup_min_ms,lookup_max_ms,lookup_ns_per_key,verify_ms," + "vram_buffers_bytes,vram_total_bytes,device_ordinal,device_name\n"; +} + +void print_csv_row(result_row const& result) +{ + std::cout << std::quoted(result.key_source) << ',' + << result.build_key_count << ',' + << result.build_key_bytes << ',' + << std::quoted(result.probe_source) << ',' + << result.probe_key_count << ',' + << result.probe_key_bytes << ',' + << result.pair_bytes << ',' + << result.output_bytes << ',' + << result.slot_bytes << ',' + << std::fixed << std::setprecision(3) + << result.requested_load_factor << ',' + << result.capacity << ',' + << result.size << ',' + << result.actual_load_factor << ',' + << result.h2d_keys_ms << ',' + << result.pair_prepare_ms << ',' + << result.insert_ms << ',' + << result.build_ms << ',' + << result.lookup_avg_ms << ',' + << result.lookup_min_ms << ',' + << result.lookup_max_ms << ',' + << result.lookup_ns_per_key << ',' + << result.verify_ms << ',' + << result.vram_buffers_bytes << ',' + << result.vram_total_bytes << ',' + << result.device_ordinal << ',' + << std::quoted(result.device_name) << '\n'; +} + +void print_human_summary(result_row const& result) +{ + std::cout << "Build: " << result.build_key_count << " from " << result.key_source << "\n"; + std::cout << "Probe: " << result.probe_key_count << " from " << result.probe_source << "\n"; + std::cout << "Device: " << result.device_ordinal << " (" << result.device_name << ")\n"; + std::cout << std::fixed << std::setprecision(3) + << "Map: requested load factor=" << result.requested_load_factor + << ", capacity=" << result.capacity + << ", size=" << result.size + << ", actual load factor=" << result.actual_load_factor << "\n"; + std::cout << "Memory: build_keys=" << result.build_key_bytes + << " bytes, probe_keys=" << result.probe_key_bytes + << " bytes, pairs=" << result.pair_bytes + << " bytes, output=" << result.output_bytes + << " bytes, slots=" << result.slot_bytes + << " bytes, vram_buffers=" << result.vram_buffers_bytes + << " bytes, vram_total=" << result.vram_total_bytes << " bytes\n"; + std::cout << "Timings (ms): h2d_keys=" << result.h2d_keys_ms + << ", pair_prepare=" << result.pair_prepare_ms + << ", insert=" << result.insert_ms + << ", build=" << result.build_ms + << ", lookup_avg=" << result.lookup_avg_ms + << ", lookup_min=" << result.lookup_min_ms + << ", lookup_max=" << result.lookup_max_ms + << ", verify=" << result.verify_ms << "\n"; + std::cout << "Throughput: " << result.lookup_ns_per_key << " ns/key\n"; +} + +} // namespace + +int main(int argc, char** argv) +{ + try { + auto const opts = parse_args(argc, argv); + + throw_if_cuda(cudaSetDevice(opts.device_ordinal), "cudaSetDevice"); + + cudaDeviceProp prop{}; + throw_if_cuda(cudaGetDeviceProperties(&prop, opts.device_ordinal), "cudaGetDeviceProperties"); + + auto const host_build_keys = load_keys_file(opts.keys_file, opts.max_keys); + auto const build_key_count = host_build_keys.size(); + if (build_key_count == 0) { + throw std::runtime_error("No build keys available to benchmark"); + } + auto sorted_build_keys = host_build_keys; + std::sort(sorted_build_keys.begin(), sorted_build_keys.end()); + if (std::adjacent_find(sorted_build_keys.begin(), sorted_build_keys.end()) != + sorted_build_keys.end()) { + throw std::runtime_error("Build keys must be unique for the static_map benchmark"); + } + auto host_probe_keys = opts.probe_keys_file.empty() + ? host_build_keys + : load_keys_file(opts.probe_keys_file, opts.max_probe_keys); + if (opts.probe_keys_file.empty() && opts.max_probe_keys > 0 && + host_probe_keys.size() > opts.max_probe_keys) { + host_probe_keys.resize(static_cast(opts.max_probe_keys)); + } + auto const probe_key_count = host_probe_keys.size(); + + if (build_key_count > static_cast(std::numeric_limits::max())) { + throw std::runtime_error("Key count exceeds uint32_t ordinal range"); + } + + std::size_t free_before_buffers = 0; + std::size_t total_device_mem = 0; + throw_if_cuda(cudaMemGetInfo(&free_before_buffers, &total_device_mem), "cudaMemGetInfo(before buffers)"); + + thrust::device_vector d_build_keys(build_key_count); + thrust::device_vector d_pairs(build_key_count); + thrust::device_vector d_probe_keys(probe_key_count); + thrust::device_vector d_values(probe_key_count); + + std::size_t free_after_buffers = 0; + throw_if_cuda(cudaMemGetInfo(&free_after_buffers, &total_device_mem), "cudaMemGetInfo(after buffers)"); + + auto constexpr empty_key = std::numeric_limits::max(); + auto constexpr empty_value = std::numeric_limits::max(); + if (std::find(host_build_keys.begin(), host_build_keys.end(), empty_key) != host_build_keys.end()) { + throw std::runtime_error("Build keys contain the reserved empty-key sentinel"); + } + if (std::find(host_probe_keys.begin(), host_probe_keys.end(), empty_key) != host_probe_keys.end()) { + throw std::runtime_error("Probe keys contain the reserved empty-key sentinel"); + } + + auto map = cuco::static_map{build_key_count, + opts.load_factor, + cuco::empty_key{empty_key}, + cuco::empty_value{empty_value}}; + + std::size_t free_after_map = 0; + throw_if_cuda(cudaMemGetInfo(&free_after_map, &total_device_mem), "cudaMemGetInfo(after map)"); + + cuda_stream_handle stream{}; + throw_if_cuda(cudaStreamCreate(&stream.value), "cudaStreamCreate"); + cuda_event_handle start{}; + cuda_event_handle stop{}; + throw_if_cuda(cudaEventCreate(&start.value), "cudaEventCreate(start)"); + throw_if_cuda(cudaEventCreate(&stop.value), "cudaEventCreate(stop)"); + + auto const build_blocks = static_cast((build_key_count + 255) / 256); + + throw_if_cuda(cudaEventRecord(start.value, stream.value), "cudaEventRecord(h2d start)"); + throw_if_cuda(cudaMemcpyAsync(thrust::raw_pointer_cast(d_build_keys.data()), + host_build_keys.data(), + build_key_count * sizeof(key_type), + cudaMemcpyHostToDevice, + stream.value), + "cudaMemcpyAsync(build keys)"); + throw_if_cuda(cudaEventRecord(stop.value, stream.value), "cudaEventRecord(h2d stop)"); + throw_if_cuda(cudaEventSynchronize(stop.value), "cudaEventSynchronize(h2d)"); + auto const h2d_keys_ms = ms_between(start.value, stop.value); + + throw_if_cuda(cudaEventRecord(start.value, stream.value), "cudaEventRecord(pairs start)"); + fill_pairs_kernel<<>>( + thrust::raw_pointer_cast(d_build_keys.data()), + thrust::raw_pointer_cast(d_pairs.data()), + build_key_count); + throw_if_cuda(cudaGetLastError(), "fill_pairs_kernel launch"); + throw_if_cuda(cudaEventRecord(stop.value, stream.value), "cudaEventRecord(pairs stop)"); + throw_if_cuda(cudaEventSynchronize(stop.value), "cudaEventSynchronize(pairs)"); + auto const pair_prepare_ms = ms_between(start.value, stop.value); + + throw_if_cuda(cudaEventRecord(start.value, stream.value), "cudaEventRecord(insert start)"); + map.insert_async(d_pairs.begin(), d_pairs.end(), cuda::stream_ref{stream.value}); + throw_if_cuda(cudaEventRecord(stop.value, stream.value), "cudaEventRecord(insert stop)"); + throw_if_cuda(cudaEventSynchronize(stop.value), "cudaEventSynchronize(insert)"); + auto const insert_ms = ms_between(start.value, stop.value); + auto const build_ms = h2d_keys_ms + pair_prepare_ms + insert_ms; + + throw_if_cuda(cudaMemcpyAsync(thrust::raw_pointer_cast(d_probe_keys.data()), + host_probe_keys.data(), + probe_key_count * sizeof(key_type), + cudaMemcpyHostToDevice, + stream.value), + "cudaMemcpyAsync(probe keys)"); + throw_if_cuda(cudaStreamSynchronize(stream.value), "cudaStreamSynchronize(probe copy)"); + + for (int i = 0; i < opts.warmup; ++i) { + map.find_async(d_probe_keys.begin(), d_probe_keys.end(), d_values.begin(), cuda::stream_ref{stream.value}); + throw_if_cuda(cudaStreamSynchronize(stream.value), "cudaStreamSynchronize(warmup)"); + } + + std::vector lookup_timings; + lookup_timings.reserve(static_cast(opts.iterations)); + for (int i = 0; i < opts.iterations; ++i) { + throw_if_cuda(cudaEventRecord(start.value, stream.value), "cudaEventRecord(find start)"); + map.find_async(d_probe_keys.begin(), d_probe_keys.end(), d_values.begin(), cuda::stream_ref{stream.value}); + throw_if_cuda(cudaEventRecord(stop.value, stream.value), "cudaEventRecord(find stop)"); + throw_if_cuda(cudaEventSynchronize(stop.value), "cudaEventSynchronize(find)"); + lookup_timings.push_back(ms_between(start.value, stop.value)); + } + + auto const [min_it, max_it] = std::minmax_element(lookup_timings.begin(), lookup_timings.end()); + double lookup_total = 0.0; + for (double t : lookup_timings) { lookup_total += t; } + auto const lookup_avg_ms = lookup_total / static_cast(lookup_timings.size()); + + double verify_ms = 0.0; + if (opts.verify) { + auto const verify_start = std::chrono::steady_clock::now(); + thrust::host_vector h_values = d_values; + for (std::size_t i = 0; i < probe_key_count; ++i) { + bool const in_build = + std::binary_search(sorted_build_keys.begin(), sorted_build_keys.end(), host_probe_keys[i]); + if (in_build) { + if (h_values[i] == empty_value) { + throw std::runtime_error("Verification failed at index " + std::to_string(i)); + } + if (static_cast(h_values[i]) >= build_key_count || + host_build_keys[static_cast(h_values[i])] != host_probe_keys[i]) { + throw std::runtime_error("Verification failed at index " + std::to_string(i)); + } + } else if (h_values[i] != empty_value) { + throw std::runtime_error("Verification failed at index " + std::to_string(i)); + } + } + auto const verify_stop = std::chrono::steady_clock::now(); + verify_ms = std::chrono::duration(verify_stop - verify_start).count(); + } + + result_row result; + result.key_source = opts.keys_file; + result.probe_source = opts.probe_keys_file.empty() ? opts.keys_file : opts.probe_keys_file; + result.device_name = prop.name; + result.build_key_count = build_key_count; + result.build_key_bytes = build_key_count * sizeof(key_type); + result.probe_key_count = probe_key_count; + result.probe_key_bytes = probe_key_count * sizeof(key_type); + result.pair_bytes = build_key_count * sizeof(pair_type); + result.output_bytes = probe_key_count * sizeof(value_type); + result.capacity = map.capacity(); + result.size = map.size(cuda::stream_ref{stream.value}); + result.slot_bytes = result.capacity * sizeof(pair_type); + result.requested_load_factor = opts.load_factor; + result.actual_load_factor = + static_cast(result.size) / static_cast(result.capacity); + result.h2d_keys_ms = h2d_keys_ms; + result.pair_prepare_ms = pair_prepare_ms; + result.insert_ms = insert_ms; + result.build_ms = build_ms; + result.lookup_avg_ms = lookup_avg_ms; + result.lookup_min_ms = *min_it; + result.lookup_max_ms = *max_it; + result.lookup_ns_per_key = ns_per_key(lookup_avg_ms, probe_key_count); + result.verify_ms = verify_ms; + result.vram_buffers_bytes = + (free_before_buffers >= free_after_buffers) ? (free_before_buffers - free_after_buffers) : 0; + result.vram_total_bytes = + (free_before_buffers >= free_after_map) ? (free_before_buffers - free_after_map) : 0; + result.device_ordinal = opts.device_ordinal; + + if (opts.csv_header) { print_csv_header(); } + if (opts.csv) { print_csv_row(result); } + else { print_human_summary(result); } + return 0; + } catch (std::exception const& e) { + std::cerr << e.what() << "\n"; + return 1; + } +} diff --git a/examples/cpp-console-cuco-static-multiset-bench/CMakeLists.txt b/examples/cpp-console-cuco-static-multiset-bench/CMakeLists.txt new file mode 100644 index 00000000..697f0c87 --- /dev/null +++ b/examples/cpp-console-cuco-static-multiset-bench/CMakeLists.txt @@ -0,0 +1,56 @@ +cmake_minimum_required(VERSION 3.20) + +project(cpp-console-cuco-static-multiset-bench LANGUAGES CXX CUDA) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CUDA_STANDARD 17) +set(CMAKE_CUDA_STANDARD_REQUIRED ON) +set(CMAKE_CUDA_EXTENSIONS OFF) + +set( + CUCO_ROOT + "" + CACHE PATH + "Path to a local cuCollections checkout." +) + +find_package(CUDAToolkit REQUIRED) + +find_path( + CUCO_INCLUDE_DIR + NAMES cuco/static_multiset.cuh + HINTS ${CUCO_ROOT} + PATH_SUFFIXES include + REQUIRED +) + +add_executable(cpp-console-cuco-static-multiset-bench + src/main.cu +) + +target_include_directories( + cpp-console-cuco-static-multiset-bench + PRIVATE + "${CUCO_INCLUDE_DIR}" +) + +target_link_libraries( + cpp-console-cuco-static-multiset-bench + PRIVATE + CUDA::cudart +) + +target_compile_options( + cpp-console-cuco-static-multiset-bench + PRIVATE + $<$:--expt-extended-lambda> +) + +if(APPLE OR UNIX) + set_property( + TARGET cpp-console-cuco-static-multiset-bench + PROPERTY BUILD_RPATH "$" + ) +endif() diff --git a/examples/cpp-console-cuco-static-multiset-bench/README.md b/examples/cpp-console-cuco-static-multiset-bench/README.md new file mode 100644 index 00000000..4002534c --- /dev/null +++ b/examples/cpp-console-cuco-static-multiset-bench/README.md @@ -0,0 +1,32 @@ +# cpp-console-cuco-static-multiset-bench + +Standalone `cuCollections` `static_multiset` baseline for the same raw `uint64_t` +`.keys` domains used by the PerfectHash GPU experiments. + +This benchmark: + +1. loads a `.keys` file as `uint64_t` keys, +2. builds a `cuco::static_multiset`, +3. runs `find_async()` over the same key stream, +4. reports build and lookup timings, including normalized `ns/key`. + +This is a closer `cuco` baseline to the join-like duplicate-preserving substrate +used in `cudf` than the `static_map` benchmark. + +## Build + +```bash +cmake -S examples/cpp-console-cuco-static-multiset-bench \ + -B build/examples/cpp-console-cuco-static-multiset-bench \ + -DCUCO_ROOT=/path/to/cucollections + +cmake --build build/examples/cpp-console-cuco-static-multiset-bench -j +``` + +## Run + +```bash +./build/examples/cpp-console-cuco-static-multiset-bench/cpp-console-cuco-static-multiset-bench \ + --keys-file /mnt/data/tpch/scale-10/perfecthash-keys/c_custkey_q03_ph-64.keys \ + --load-factor 0.5 +``` diff --git a/examples/cpp-console-cuco-static-multiset-bench/src/main.cu b/examples/cpp-console-cuco-static-multiset-bench/src/main.cu new file mode 100644 index 00000000..18ee94e4 --- /dev/null +++ b/examples/cpp-console-cuco-static-multiset-bench/src/main.cu @@ -0,0 +1,424 @@ +#include + +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +using key_type = std::uint64_t; + +struct options { + std::string keys_file; + std::string probe_keys_file; + std::uint64_t max_keys = 0; + std::uint64_t max_probe_keys = 0; + double load_factor = 0.5; + int device_ordinal = 0; + int warmup = 2; + int iterations = 10; + bool csv = false; + bool csv_header = false; + bool verify = true; +}; + +struct result_row { + std::string key_source; + std::string probe_source; + std::string device_name; + std::size_t build_key_count = 0; + std::size_t build_key_bytes = 0; + std::size_t probe_key_count = 0; + std::size_t probe_key_bytes = 0; + std::size_t output_bytes = 0; + std::size_t slot_bytes = 0; + std::size_t capacity = 0; + std::size_t size = 0; + double requested_load_factor = 0.0; + double actual_load_factor = 0.0; + double h2d_keys_ms = 0.0; + double insert_ms = 0.0; + double build_ms = 0.0; + double lookup_avg_ms = 0.0; + double lookup_min_ms = 0.0; + double lookup_max_ms = 0.0; + double lookup_ns_per_key = 0.0; + double verify_ms = 0.0; + std::size_t vram_buffers_bytes = 0; + std::size_t vram_total_bytes = 0; + int device_ordinal = 0; +}; + +struct cuda_stream_handle { + cudaStream_t value{}; + ~cuda_stream_handle() + { + if (value != nullptr) { cudaStreamDestroy(value); } + } +}; + +struct cuda_event_handle { + cudaEvent_t value{}; + ~cuda_event_handle() + { + if (value != nullptr) { cudaEventDestroy(value); } + } +}; + +void throw_if_cuda(cudaError_t status, char const* what) +{ + if (status != cudaSuccess) { + throw std::runtime_error(std::string(what) + " failed: " + cudaGetErrorString(status)); + } +} + +double ms_between(cudaEvent_t start, cudaEvent_t stop) +{ + float ms = 0.0f; + throw_if_cuda(cudaEventElapsedTime(&ms, start, stop), "cudaEventElapsedTime"); + return static_cast(ms); +} + +double ns_per_key(double milliseconds, std::size_t key_count) +{ + if (key_count == 0) { return 0.0; } + return (milliseconds * 1'000'000.0) / static_cast(key_count); +} + +std::vector load_keys_file(std::string const& path, std::uint64_t max_keys) +{ + std::ifstream input(path, std::ios::binary | std::ios::ate); + if (!input) { throw std::runtime_error("Unable to open keys file: " + path); } + + auto const size = input.tellg(); + if (size < 0) { throw std::runtime_error("Unable to determine keys file size: " + path); } + if ((static_cast(size) % sizeof(key_type)) != 0) { + throw std::runtime_error("Keys file size is not a multiple of 8 bytes: " + path); + } + + auto const available_keys = + static_cast(size) / static_cast(sizeof(key_type)); + auto const keys_to_read = + (max_keys != 0 && max_keys < available_keys) ? max_keys : available_keys; + + std::vector keys(static_cast(keys_to_read)); + input.seekg(0, std::ios::beg); + input.read(reinterpret_cast(keys.data()), + static_cast(keys.size() * sizeof(key_type))); + if (!input) { throw std::runtime_error("Unable to read keys file contents: " + path); } + + return keys; +} + +void print_usage(char const* argv0) +{ + std::cout << "Usage: " << argv0 + << " --keys-file [--probe-keys-file ] [--max-keys ] [--max-probe-keys ] [--load-factor ]\n" + " [--device ] [--warmup ] [--iterations ]\n" + " [--csv] [--csv-header] [--no-verify]\n"; +} + +options parse_args(int argc, char** argv) +{ + options opts; + for (int index = 1; index < argc; ++index) { + std::string arg = argv[index]; + if (arg == "--help" || arg == "-h") { + print_usage(argv[0]); + std::exit(0); + } + if (arg == "--keys-file" && index + 1 < argc) { + opts.keys_file = argv[++index]; + continue; + } + if (arg == "--probe-keys-file" && index + 1 < argc) { + opts.probe_keys_file = argv[++index]; + continue; + } + if (arg == "--max-keys" && index + 1 < argc) { + opts.max_keys = std::stoull(argv[++index]); + continue; + } + if (arg == "--max-probe-keys" && index + 1 < argc) { + opts.max_probe_keys = std::stoull(argv[++index]); + continue; + } + if (arg == "--load-factor" && index + 1 < argc) { + opts.load_factor = std::stod(argv[++index]); + continue; + } + if (arg == "--device" && index + 1 < argc) { + opts.device_ordinal = std::stoi(argv[++index]); + continue; + } + if (arg == "--warmup" && index + 1 < argc) { + opts.warmup = std::stoi(argv[++index]); + continue; + } + if (arg == "--iterations" && index + 1 < argc) { + opts.iterations = std::stoi(argv[++index]); + continue; + } + if (arg == "--csv") { + opts.csv = true; + continue; + } + if (arg == "--csv-header") { + opts.csv = true; + opts.csv_header = true; + continue; + } + if (arg == "--no-verify") { + opts.verify = false; + continue; + } + throw std::runtime_error("Unknown argument: " + arg); + } + + if (opts.keys_file.empty()) { throw std::runtime_error("--keys-file is required"); } + if (!(opts.load_factor > 0.0 && opts.load_factor < 1.0)) { + throw std::runtime_error("--load-factor must be in the open interval (0, 1)"); + } + if (opts.warmup < 0) { throw std::runtime_error("--warmup must be >= 0"); } + if (opts.iterations <= 0) { throw std::runtime_error("--iterations must be > 0"); } + + return opts; +} + +void print_csv_header() +{ + std::cout << "key_source,build_key_count,build_key_bytes,probe_source,probe_key_count,probe_key_bytes,output_bytes,slot_bytes,requested_load_factor," + "capacity,size,actual_load_factor,h2d_keys_ms,insert_ms,build_ms," + "lookup_avg_ms,lookup_min_ms,lookup_max_ms,lookup_ns_per_key,verify_ms," + "vram_buffers_bytes,vram_total_bytes,device_ordinal,device_name\n"; +} + +void print_csv_row(result_row const& result) +{ + std::cout << std::quoted(result.key_source) << ',' + << result.build_key_count << ',' + << result.build_key_bytes << ',' + << std::quoted(result.probe_source) << ',' + << result.probe_key_count << ',' + << result.probe_key_bytes << ',' + << result.output_bytes << ',' + << result.slot_bytes << ',' + << std::fixed << std::setprecision(3) + << result.requested_load_factor << ',' + << result.capacity << ',' + << result.size << ',' + << result.actual_load_factor << ',' + << result.h2d_keys_ms << ',' + << result.insert_ms << ',' + << result.build_ms << ',' + << result.lookup_avg_ms << ',' + << result.lookup_min_ms << ',' + << result.lookup_max_ms << ',' + << result.lookup_ns_per_key << ',' + << result.verify_ms << ',' + << result.vram_buffers_bytes << ',' + << result.vram_total_bytes << ',' + << result.device_ordinal << ',' + << std::quoted(result.device_name) << '\n'; +} + +void print_human_summary(result_row const& result) +{ + std::cout << "Build: " << result.build_key_count << " from " << result.key_source << "\n"; + std::cout << "Probe: " << result.probe_key_count << " from " << result.probe_source << "\n"; + std::cout << "Device: " << result.device_ordinal << " (" << result.device_name << ")\n"; + std::cout << std::fixed << std::setprecision(3) + << "Set: requested load factor=" << result.requested_load_factor + << ", capacity=" << result.capacity + << ", size=" << result.size + << ", actual load factor=" << result.actual_load_factor << "\n"; + std::cout << "Memory: build_keys=" << result.build_key_bytes + << " bytes, probe_keys=" << result.probe_key_bytes + << " bytes, output=" << result.output_bytes + << " bytes, slots=" << result.slot_bytes + << " bytes, vram_buffers=" << result.vram_buffers_bytes + << " bytes, vram_total=" << result.vram_total_bytes << " bytes\n"; + std::cout << "Timings (ms): h2d_keys=" << result.h2d_keys_ms + << ", insert=" << result.insert_ms + << ", build=" << result.build_ms + << ", lookup_avg=" << result.lookup_avg_ms + << ", lookup_min=" << result.lookup_min_ms + << ", lookup_max=" << result.lookup_max_ms + << ", verify=" << result.verify_ms << "\n"; + std::cout << "Throughput: " << result.lookup_ns_per_key << " ns/key\n"; +} + +} // namespace + +int main(int argc, char** argv) +{ + try { + auto const opts = parse_args(argc, argv); + + throw_if_cuda(cudaSetDevice(opts.device_ordinal), "cudaSetDevice"); + + cudaDeviceProp prop{}; + throw_if_cuda(cudaGetDeviceProperties(&prop, opts.device_ordinal), "cudaGetDeviceProperties"); + + auto const host_build_keys = load_keys_file(opts.keys_file, opts.max_keys); + auto const build_key_count = host_build_keys.size(); + if (build_key_count == 0) { + throw std::runtime_error("No build keys available to benchmark"); + } + auto sorted_build_keys = host_build_keys; + std::sort(sorted_build_keys.begin(), sorted_build_keys.end()); + auto host_probe_keys = opts.probe_keys_file.empty() + ? host_build_keys + : load_keys_file(opts.probe_keys_file, opts.max_probe_keys); + if (opts.probe_keys_file.empty() && opts.max_probe_keys > 0 && + host_probe_keys.size() > opts.max_probe_keys) { + host_probe_keys.resize(static_cast(opts.max_probe_keys)); + } + auto const probe_key_count = host_probe_keys.size(); + + std::size_t free_before_buffers = 0; + std::size_t total_device_mem = 0; + throw_if_cuda(cudaMemGetInfo(&free_before_buffers, &total_device_mem), "cudaMemGetInfo(before buffers)"); + + thrust::device_vector d_build_keys(build_key_count); + thrust::device_vector d_probe_keys(probe_key_count); + thrust::device_vector d_output(probe_key_count); + + std::size_t free_after_buffers = 0; + throw_if_cuda(cudaMemGetInfo(&free_after_buffers, &total_device_mem), "cudaMemGetInfo(after buffers)"); + + auto constexpr empty_key = std::numeric_limits::max(); + if (std::find(host_build_keys.begin(), host_build_keys.end(), empty_key) != host_build_keys.end()) { + throw std::runtime_error("Build keys contain the reserved empty-key sentinel"); + } + if (std::find(host_probe_keys.begin(), host_probe_keys.end(), empty_key) != host_probe_keys.end()) { + throw std::runtime_error("Probe keys contain the reserved empty-key sentinel"); + } + + auto set = cuco::static_multiset{build_key_count, opts.load_factor, cuco::empty_key{empty_key}}; + + std::size_t free_after_set = 0; + throw_if_cuda(cudaMemGetInfo(&free_after_set, &total_device_mem), "cudaMemGetInfo(after set)"); + + cuda_stream_handle stream{}; + throw_if_cuda(cudaStreamCreate(&stream.value), "cudaStreamCreate"); + cuda_event_handle start{}; + cuda_event_handle stop{}; + throw_if_cuda(cudaEventCreate(&start.value), "cudaEventCreate(start)"); + throw_if_cuda(cudaEventCreate(&stop.value), "cudaEventCreate(stop)"); + + throw_if_cuda(cudaEventRecord(start.value, stream.value), "cudaEventRecord(h2d start)"); + throw_if_cuda(cudaMemcpyAsync(thrust::raw_pointer_cast(d_build_keys.data()), + host_build_keys.data(), + build_key_count * sizeof(key_type), + cudaMemcpyHostToDevice, + stream.value), + "cudaMemcpyAsync(build keys)"); + throw_if_cuda(cudaEventRecord(stop.value, stream.value), "cudaEventRecord(h2d stop)"); + throw_if_cuda(cudaEventSynchronize(stop.value), "cudaEventSynchronize(h2d)"); + auto const h2d_keys_ms = ms_between(start.value, stop.value); + + throw_if_cuda(cudaEventRecord(start.value, stream.value), "cudaEventRecord(insert start)"); + set.insert_async(d_build_keys.begin(), d_build_keys.end(), cuda::stream_ref{stream.value}); + throw_if_cuda(cudaEventRecord(stop.value, stream.value), "cudaEventRecord(insert stop)"); + throw_if_cuda(cudaEventSynchronize(stop.value), "cudaEventSynchronize(insert)"); + auto const insert_ms = ms_between(start.value, stop.value); + auto const build_ms = h2d_keys_ms + insert_ms; + + throw_if_cuda(cudaMemcpyAsync(thrust::raw_pointer_cast(d_probe_keys.data()), + host_probe_keys.data(), + probe_key_count * sizeof(key_type), + cudaMemcpyHostToDevice, + stream.value), + "cudaMemcpyAsync(probe keys)"); + throw_if_cuda(cudaStreamSynchronize(stream.value), "cudaStreamSynchronize(probe copy)"); + + for (int i = 0; i < opts.warmup; ++i) { + set.find_async(d_probe_keys.begin(), d_probe_keys.end(), d_output.begin(), cuda::stream_ref{stream.value}); + throw_if_cuda(cudaStreamSynchronize(stream.value), "cudaStreamSynchronize(warmup)"); + } + + std::vector lookup_timings; + lookup_timings.reserve(static_cast(opts.iterations)); + for (int i = 0; i < opts.iterations; ++i) { + throw_if_cuda(cudaEventRecord(start.value, stream.value), "cudaEventRecord(find start)"); + set.find_async(d_probe_keys.begin(), d_probe_keys.end(), d_output.begin(), cuda::stream_ref{stream.value}); + throw_if_cuda(cudaEventRecord(stop.value, stream.value), "cudaEventRecord(find stop)"); + throw_if_cuda(cudaEventSynchronize(stop.value), "cudaEventSynchronize(find)"); + lookup_timings.push_back(ms_between(start.value, stop.value)); + } + + auto const [min_it, max_it] = std::minmax_element(lookup_timings.begin(), lookup_timings.end()); + double lookup_total = 0.0; + for (double t : lookup_timings) { lookup_total += t; } + auto const lookup_avg_ms = lookup_total / static_cast(lookup_timings.size()); + + double verify_ms = 0.0; + if (opts.verify) { + auto const verify_start = std::chrono::steady_clock::now(); + thrust::host_vector h_output = d_output; + for (std::size_t i = 0; i < probe_key_count; ++i) { + bool const in_build = + std::binary_search(sorted_build_keys.begin(), sorted_build_keys.end(), host_probe_keys[i]); + if (in_build) { + if (h_output[i] == empty_key || h_output[i] != host_probe_keys[i]) { + throw std::runtime_error("Verification failed at index " + std::to_string(i)); + } + } else if (h_output[i] != empty_key) { + throw std::runtime_error("Verification failed at index " + std::to_string(i)); + } + } + auto const verify_stop = std::chrono::steady_clock::now(); + verify_ms = std::chrono::duration(verify_stop - verify_start).count(); + } + + result_row result; + result.key_source = opts.keys_file; + result.probe_source = opts.probe_keys_file.empty() ? opts.keys_file : opts.probe_keys_file; + result.device_name = prop.name; + result.build_key_count = build_key_count; + result.build_key_bytes = build_key_count * sizeof(key_type); + result.probe_key_count = probe_key_count; + result.probe_key_bytes = probe_key_count * sizeof(key_type); + result.output_bytes = probe_key_count * sizeof(key_type); + result.capacity = set.capacity(); + result.size = set.size(cuda::stream_ref{stream.value}); + result.slot_bytes = result.capacity * sizeof(key_type); + result.requested_load_factor = opts.load_factor; + result.actual_load_factor = + static_cast(result.size) / static_cast(result.capacity); + result.h2d_keys_ms = h2d_keys_ms; + result.insert_ms = insert_ms; + result.build_ms = build_ms; + result.lookup_avg_ms = lookup_avg_ms; + result.lookup_min_ms = *min_it; + result.lookup_max_ms = *max_it; + result.lookup_ns_per_key = ns_per_key(lookup_avg_ms, probe_key_count); + result.verify_ms = verify_ms; + result.vram_buffers_bytes = + (free_before_buffers >= free_after_buffers) ? (free_before_buffers - free_after_buffers) : 0; + result.vram_total_bytes = + (free_before_buffers >= free_after_set) ? (free_before_buffers - free_after_set) : 0; + result.device_ordinal = opts.device_ordinal; + + if (opts.csv_header) { print_csv_header(); } + if (opts.csv) { print_csv_row(result); } + else { print_human_summary(result); } + return 0; + } catch (std::exception const& e) { + std::cerr << e.what() << "\n"; + return 1; + } +} diff --git a/examples/cpp-console-online-cuda-nvrtc/CMakeLists.txt b/examples/cpp-console-online-cuda-nvrtc/CMakeLists.txt new file mode 100644 index 00000000..33ad9fac --- /dev/null +++ b/examples/cpp-console-online-cuda-nvrtc/CMakeLists.txt @@ -0,0 +1,210 @@ +cmake_minimum_required(VERSION 3.20) + +project(cpp-console-online-cuda-nvrtc LANGUAGES CXX) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) + +option( + PH_ONLINE_CUDA_NVRTC_PREFER_SHARED + "Prefer shared PerfectHash runtime over static." + ON +) +option( + PH_ONLINE_CUDA_NVRTC_USE_FETCHCONTENT + "Fetch PerfectHash from GitHub via FetchContent." + ON +) + +set( + PERFECTHASH_GIT_REPOSITORY + "https://github.com/tpn/perfecthash.git" + CACHE STRING + "PerfectHash git repository used with FetchContent." +) +set( + PERFECTHASH_GIT_TAG + "main" + CACHE STRING + "PerfectHash git ref used with FetchContent." +) +set( + PERFECTHASH_BUILD_PROFILE + "online-rawdog-jit" + CACHE STRING + "PerfectHash profile used with FetchContent." +) +set( + PERFECTHASH_ROOT + "" + CACHE PATH + "Path hint to a PerfectHash repo root, build tree, or install prefix." +) + +find_package(CUDAToolkit REQUIRED) +list(GET CUDAToolkit_INCLUDE_DIRS 0 PH_ONLINE_CUDA_NVRTC_CUDA_INCLUDE_DIR) + +set(PH_ONLINE_CUDA_NVRTC_TARGET "") +set(PH_ONLINE_CUDA_NVRTC_RUNTIME_TARGET "") +set(PH_ONLINE_CUDA_NVRTC_LLVM_TARGET "") + +if(PH_ONLINE_CUDA_NVRTC_USE_FETCHCONTENT) + include(FetchContent) + + set(PERFECTHASH_ENABLE_TESTS OFF CACHE BOOL "" FORCE) + set(BUILD_TESTING OFF CACHE BOOL "" FORCE) + set(PERFECTHASH_ENABLE_INSTALL ON CACHE BOOL "" FORCE) + set(PERFECTHASH_BUILD_PROFILE "${PERFECTHASH_BUILD_PROFILE}" CACHE STRING "" FORCE) + + FetchContent_Declare( + PerfectHash + GIT_REPOSITORY ${PERFECTHASH_GIT_REPOSITORY} + GIT_TAG ${PERFECTHASH_GIT_TAG} + ) + FetchContent_MakeAvailable(PerfectHash) + + if(TARGET PerfectHash::PerfectHash) + set(PH_ONLINE_CUDA_NVRTC_TARGET PerfectHash::PerfectHash) + set(PH_ONLINE_CUDA_NVRTC_RUNTIME_TARGET PerfectHash::PerfectHash) + elseif(TARGET PerfectHash::PerfectHashOnlineCore) + set(PH_ONLINE_CUDA_NVRTC_TARGET PerfectHash::PerfectHashOnlineCore) + set(PH_ONLINE_CUDA_NVRTC_RUNTIME_TARGET PerfectHash::PerfectHashOnlineCore) + elseif(TARGET PerfectHash::PerfectHashOnline) + set(PH_ONLINE_CUDA_NVRTC_TARGET PerfectHash::PerfectHashOnline) + set(PH_ONLINE_CUDA_NVRTC_RUNTIME_TARGET PerfectHash::PerfectHashOnline) + else() + message( + FATAL_ERROR + "FetchContent build did not expose PerfectHash online targets. " + "Verify PERFECTHASH_BUILD_PROFILE='${PERFECTHASH_BUILD_PROFILE}'." + ) + endif() + + if(TARGET PerfectHash::PerfectHashLLVM) + set(PH_ONLINE_CUDA_NVRTC_LLVM_TARGET PerfectHash::PerfectHashLLVM) + endif() +else() + list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../cpp-console-online-jit/cmake") + set(PH_ONLINE_JIT_PREFER_SHARED "${PH_ONLINE_CUDA_NVRTC_PREFER_SHARED}") + find_package(PerfectHashOnlineJit REQUIRED) + set(PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY "") + if(EXISTS "${PERFECTHASH_ROOT}/build-verify-full/lib/libPerfectHash.so") + set(PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY + "${PERFECTHASH_ROOT}/build-verify-full/lib/libPerfectHash.so") + elseif(EXISTS "${PERFECTHASH_ROOT}/lib/libPerfectHash.so") + set(PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY + "${PERFECTHASH_ROOT}/lib/libPerfectHash.so") + else() + find_library( + PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY + NAMES PerfectHash + HINTS ${PERFECTHASH_ROOT} + PATH_SUFFIXES + lib + build/lib + build-verify-full/lib + ) + endif() + if(PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY) + set(PH_ONLINE_CUDA_NVRTC_TARGET "${PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY}") + if((NOT DEFINED PERFECTHASH_ONLINE_JIT_INCLUDE_DIR OR + PERFECTHASH_ONLINE_JIT_INCLUDE_DIR STREQUAL "") AND + EXISTS "${PERFECTHASH_ROOT}/include/PerfectHash/PerfectHashOnlineJit.h") + set(PERFECTHASH_ONLINE_JIT_INCLUDE_DIR "${PERFECTHASH_ROOT}/include") + endif() + else() + set(PH_ONLINE_CUDA_NVRTC_TARGET PerfectHash::OnlineJit) + endif() +endif() + +add_executable(cpp-console-online-cuda-nvrtc + src/main.cpp +) + +target_link_libraries(cpp-console-online-cuda-nvrtc + PRIVATE + ${PH_ONLINE_CUDA_NVRTC_TARGET} + CUDA::cudart + CUDA::nvrtc + CUDA::nvJitLink + CUDA::cuda_driver +) + +if(DEFINED PERFECTHASH_ONLINE_JIT_INCLUDE_DIR AND + NOT PERFECTHASH_ONLINE_JIT_INCLUDE_DIR STREQUAL "") + get_filename_component(_ph_include_name "${PERFECTHASH_ONLINE_JIT_INCLUDE_DIR}" NAME) + target_include_directories( + cpp-console-online-cuda-nvrtc + PRIVATE + "${PERFECTHASH_ONLINE_JIT_INCLUDE_DIR}" + ) + if(_ph_include_name STREQUAL "PerfectHash") + get_filename_component(_ph_include_parent "${PERFECTHASH_ONLINE_JIT_INCLUDE_DIR}" DIRECTORY) + target_include_directories( + cpp-console-online-cuda-nvrtc + PRIVATE + "${_ph_include_parent}" + ) + endif() +endif() + +target_compile_definitions( + cpp-console-online-cuda-nvrtc + PRIVATE + PH_CUDA_INCLUDE_DIR="${PH_ONLINE_CUDA_NVRTC_CUDA_INCLUDE_DIR}" +) + +if(PH_ONLINE_CUDA_NVRTC_LLVM_TARGET) + target_compile_definitions( + cpp-console-online-cuda-nvrtc + PRIVATE + PH_ONLINE_JIT_LLVM_LIBRARY_PATH="$" + ) +elseif(DEFINED PERFECTHASH_ONLINE_JIT_LLVM_LIBRARY AND + NOT PERFECTHASH_ONLINE_JIT_LLVM_LIBRARY STREQUAL "") + target_compile_definitions( + cpp-console-online-cuda-nvrtc + PRIVATE + PH_ONLINE_JIT_LLVM_LIBRARY_PATH="${PERFECTHASH_ONLINE_JIT_LLVM_LIBRARY}" + ) +elseif(EXISTS "${PERFECTHASH_ROOT}/build-verify-full/lib/libPerfectHashLLVM.so") + target_compile_definitions( + cpp-console-online-cuda-nvrtc + PRIVATE + PH_ONLINE_JIT_LLVM_LIBRARY_PATH="${PERFECTHASH_ROOT}/build-verify-full/lib/libPerfectHashLLVM.so" + ) +elseif(EXISTS "${PERFECTHASH_ROOT}/lib/libPerfectHashLLVM.so") + target_compile_definitions( + cpp-console-online-cuda-nvrtc + PRIVATE + PH_ONLINE_JIT_LLVM_LIBRARY_PATH="${PERFECTHASH_ROOT}/lib/libPerfectHashLLVM.so" + ) +endif() + +set(_ph_runtime_dirs) +if(PH_ONLINE_CUDA_NVRTC_RUNTIME_TARGET) + list(APPEND _ph_runtime_dirs "$") +endif() +if(PH_ONLINE_CUDA_NVRTC_LLVM_TARGET) + list(APPEND _ph_runtime_dirs "$") +endif() +if(DEFINED PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY AND + NOT PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY STREQUAL "") + get_filename_component(_ph_full_lib_dir "${PH_ONLINE_CUDA_NVRTC_FULL_LIBRARY}" DIRECTORY) + list(APPEND _ph_runtime_dirs "${_ph_full_lib_dir}") +endif() +list(APPEND _ph_runtime_dirs + "$" + "$" + "$" +) +list(REMOVE_DUPLICATES _ph_runtime_dirs) + +if(_ph_runtime_dirs AND (APPLE OR UNIX)) + string(JOIN ";" _ph_build_rpath ${_ph_runtime_dirs}) + set_property( + TARGET cpp-console-online-cuda-nvrtc + PROPERTY BUILD_RPATH "${_ph_build_rpath}" + ) +endif() diff --git a/examples/cpp-console-online-cuda-nvrtc/README.md b/examples/cpp-console-online-cuda-nvrtc/README.md new file mode 100644 index 00000000..e81210b1 --- /dev/null +++ b/examples/cpp-console-online-cuda-nvrtc/README.md @@ -0,0 +1,158 @@ +# cpp-console-online-cuda-nvrtc + +Minimal C++ example and benchmark driver that: + +1. builds a perfect hash table on CPU via `PerfectHashOnlineJit`, +2. requests the fragment-only generated CUDA lookup source, +3. uploads the assigned table payload separately by default, +4. appends its own templated `ITEMS_PER_THREAD` probe kernel, +5. compiles the combined translation unit with NVRTC for the active GPU, and +6. launches the kernel via the CUDA driver API. + +This is the intended integration shape for downstream GPU consumers that want +call-site specialization around the lookup routine, rather than relying on a +prepackaged kernel shape. + +By default, the generated CUDA fragment omits the embedded `table_data[]` array +and expects the consumer to provide the assigned table payload as a runtime +kernel argument. This keeps the NVRTC translation unit small even for large +TPC-H domains. Use `--embed-table-data` to restore the previous fully embedded +mode for comparison. + +It can run against either: + +- the built-in synthetic 64-bit sample domain, or +- a raw little-endian `uint64_t` `.keys` file such as the extracted TPC-H + domains under `/mnt/data/tpch/.../perfecthash-keys/*.keys` + +## Build + +### Linux/macOS + +```bash +cmake -S examples/cpp-console-online-cuda-nvrtc \ + -B build/examples/cpp-console-online-cuda-nvrtc \ + -DPERFECTHASH_BUILD_PROFILE=online-rawdog-jit + +cmake --build build/examples/cpp-console-online-cuda-nvrtc -j +``` + +### Local Tree Fallback + +```bash +cmake -S examples/cpp-console-online-cuda-nvrtc \ + -B build/examples/cpp-console-online-cuda-nvrtc \ + -DPH_ONLINE_CUDA_NVRTC_USE_FETCHCONTENT=OFF \ + -DPERFECTHASH_ROOT=/path/to/perfecthash/build-verify-full +``` + +## Run + +```bash +./build/examples/cpp-console-online-cuda-nvrtc/cpp-console-online-cuda-nvrtc \ + --hash mulshrolate3rx \ + --items-per-thread 8 +``` + +Optional arguments: + +- `--keys-file ` +- `--max-keys ` +- `--hash ` +- `--items-per-thread ` +- `--threads ` +- `--iterations ` +- `--warmup ` +- `--device ` +- `--compile-mode ` +- `--lookup-mode ` +- `--table-load-mode ` +- `--cpu-backend ` +- `--cpu-vector-width <1|2|4|8|16>` +- `--cpu-strict-vector-width <0|1>` +- `--analyze-slot-reuse` +- `--analysis-only` +- `--dump-fragment` +- `--source-out ` +- `--embed-table-data` +- `--csv` +- `--csv-header` +- `--no-verify` + +`--dump-fragment` prints the fragment-only generated CUDA source. + +`--source-out` writes the full combined NVRTC translation unit to disk, which is +useful for debugging or for experimenting with PTX versus LTO-IR compilation +strategies. + +`--embed-table-data` forces the generated fragment to inline the assigned table +payload into the CUDA source instead of uploading it separately at runtime. + +`--lookup-mode direct` runs the current fused path where each item computes its +two assigned-table offsets and immediately performs the two table reads. + +`--lookup-mode split` breaks the work into two kernels: + +- `compute_slots_kernel`: compute the two assigned-table offsets per key +- `gather_kernel`: perform the assigned-table reads and final index add + +This is intended for backend analysis. It does not reorder or coalesce the +requests yet; it simply separates address generation from the random gathers. +Like the other non-`direct` modes, it is currently intended for self-probe or +known-member workloads, not external build/probe-stream validation. + +`--lookup-mode warpcache` keeps the fused structure but uses warp-local +duplicate detection for slot loads, so identical slot requests within a warp +can be serviced once and broadcast with warp intrinsics. It is currently +intended for self-probe or known-member workloads, not external build/probe +stream validation. + +`--lookup-mode blocksort` is a heavier experimental path that sorts all slot +requests for a block in shared memory before performing the assigned-table +loads, then scatters the gathered values back to the original outputs. It is a +clustering experiment, not a production path, and is currently intended for +self-probe or known-member workloads. + +`--table-load-mode generic` uses normal global loads for the assigned table. + +`--table-load-mode readonly` routes assigned-table reads through an explicit +read-only load helper in the generated consumer TU. + +`--cpu-backend` enables a CPU bulk-index baseline in the same run. The benchmark +derives exact downsized 32-bit keys from the 64-bit source domain using the +table's downsize bitmap and then runs the requested JIT backend/vector width via +the public `Index32`/`Index32xN` APIs. + +`--cpu-vector-width` selects the requested CPU bulk width. The CSV output +records both requested and effective width, so fallback behavior is visible. + +`--cpu-strict-vector-width` forwards +`PH_ONLINE_JIT_COMPILE_FLAG_STRICT_VECTOR_WIDTH` to the CPU JIT compile path. + +`--analyze-slot-reuse` computes exact host-side slot-stream reuse statistics for +the current `ITEMS_PER_THREAD` and block size using the table's actual seeds, +masks, and downsize bitmap. + +`--analysis-only` runs the slot-reuse analysis and skips the CPU/GPU benchmark +phases. + +`--csv` prints a single machine-readable result row. `--csv-header` prepends the +header line. Both CPU and GPU lookup throughput are reported as normalized +`ns/key` figures; GPU values can legitimately be sub-nanosecond because they are +aggregate throughput numbers over many concurrent threads. + +`--no-verify` skips the post-kernel uniqueness check on returned indices. + +## TPC-H Example + +```bash +./build/examples/cpp-console-online-cuda-nvrtc/cpp-console-online-cuda-nvrtc \ + --keys-file /mnt/data/tpch/scale-10/perfecthash-keys/c_custkey_q03_ph-64.keys \ + --hash mulshrolate3rx \ + --compile-mode lto \ + --lookup-mode direct \ + --items-per-thread 8 \ + --threads 128 \ + --warmup 2 \ + --iterations 10 +``` diff --git a/examples/cpp-console-online-cuda-nvrtc/src/main.cpp b/examples/cpp-console-online-cuda-nvrtc/src/main.cpp new file mode 100644 index 00000000..81dbf438 --- /dev/null +++ b/examples/cpp-console-online-cuda-nvrtc/src/main.cpp @@ -0,0 +1,2676 @@ +#include + +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#if !defined(_WIN32) && defined(PH_ONLINE_JIT_LLVM_LIBRARY_PATH) +#include +#endif + +#if !defined(_WIN32) +#include +#include +#endif + +namespace { + +using steady_clock = std::chrono::steady_clock; + +std::string to_hex(std::uint32_t value) +{ + std::ostringstream stream; + stream << "0x" << std::uppercase << std::hex << value; + return stream.str(); +} + +bool succeeded(std::int32_t hr) { return hr >= 0; } + +bool is_supported_hash_name(std::string const& name) +{ + return (name == "multiplyshiftr" || name == "multiplyshiftlr" || + name == "multiplyshiftrmultiply" || name == "multiplyshiftr2" || + name == "multiplyshiftrx" || + name == "mulshrolate1rx" || name == "mulshrolate2rx" || + name == "mulshrolate3rx" || name == "mulshrolate4rx"); +} + +PH_ONLINE_JIT_HASH_FUNCTION parse_hash_function(std::string const& name) +{ + if (name == "multiplyshiftr") { return PhOnlineJitHashMultiplyShiftR; } + if (name == "multiplyshiftlr") { return PhOnlineJitHashMultiplyShiftLR; } + if (name == "multiplyshiftrmultiply") { return PhOnlineJitHashMultiplyShiftRMultiply; } + if (name == "multiplyshiftr2") { return PhOnlineJitHashMultiplyShiftR2; } + if (name == "multiplyshiftrx") { return PhOnlineJitHashMultiplyShiftRX; } + if (name == "mulshrolate1rx") { return PhOnlineJitHashMulshrolate1RX; } + if (name == "mulshrolate2rx") { return PhOnlineJitHashMulshrolate2RX; } + if (name == "mulshrolate3rx") { return PhOnlineJitHashMulshrolate3RX; } + if (name == "mulshrolate4rx") { return PhOnlineJitHashMulshrolate4RX; } + return PhOnlineJitHashMulshrolate3RX; +} + +enum class compile_mode { ptx, lto }; +enum class lookup_mode { direct, split, warpcache, blocksort }; +enum class table_load_mode { generic, readonly }; + +compile_mode parse_compile_mode(std::string const& name) +{ + if (name == "ptx") { return compile_mode::ptx; } + if (name == "lto") { return compile_mode::lto; } + return compile_mode::ptx; +} + +char const* compile_mode_to_string(compile_mode mode) +{ + switch (mode) { + case compile_mode::ptx: return "ptx"; + case compile_mode::lto: return "lto"; + } + return "unknown"; +} + +bool is_supported_backend_name(std::string const& name) +{ + return (name == "none" || name == "auto" || name == "rawdog-jit" || + name == "llvm-jit"); +} + +PH_ONLINE_JIT_BACKEND parse_backend(std::string const& name) +{ + if (name == "rawdog-jit") { return PhOnlineJitBackendRawDogJit; } + if (name == "llvm-jit") { return PhOnlineJitBackendLlvmJit; } + if (name == "auto") { return PhOnlineJitBackendAuto; } + return PhOnlineJitBackendRawDogJit; +} + +char const* backend_to_string(PH_ONLINE_JIT_BACKEND backend) +{ + switch (backend) { + case PhOnlineJitBackendAuto: return "auto"; + case PhOnlineJitBackendRawDogJit: return "rawdog-jit"; + case PhOnlineJitBackendLlvmJit: return "llvm-jit"; + } + return "unknown"; +} + +void preload_llvm_runtime_library(std::string const& backend_name) +{ +#if !defined(_WIN32) && defined(PH_ONLINE_JIT_LLVM_LIBRARY_PATH) + if (backend_name == "llvm-jit" || backend_name == "auto") { + void* handle = dlopen(PH_ONLINE_JIT_LLVM_LIBRARY_PATH, RTLD_NOW | RTLD_GLOBAL); + if (!handle) { + std::cerr << "Warning: unable to preload LLVM runtime from " + << PH_ONLINE_JIT_LLVM_LIBRARY_PATH << ": " << dlerror() << "\n"; + } + } +#else + (void)backend_name; +#endif +} + +lookup_mode parse_lookup_mode(std::string const& name) +{ + if (name == "direct") { return lookup_mode::direct; } + if (name == "split") { return lookup_mode::split; } + if (name == "warpcache") { return lookup_mode::warpcache; } + if (name == "blocksort") { return lookup_mode::blocksort; } + return lookup_mode::direct; +} + +char const* lookup_mode_to_string(lookup_mode mode) +{ + switch (mode) { + case lookup_mode::direct: return "direct"; + case lookup_mode::split: return "split"; + case lookup_mode::warpcache: return "warpcache"; + case lookup_mode::blocksort: return "blocksort"; + } + return "unknown"; +} + +table_load_mode parse_table_load_mode(std::string const& name) +{ + if (name == "generic") { return table_load_mode::generic; } + if (name == "readonly") { return table_load_mode::readonly; } + return table_load_mode::generic; +} + +char const* table_load_mode_to_string(table_load_mode mode) +{ + switch (mode) { + case table_load_mode::generic: return "generic"; + case table_load_mode::readonly: return "readonly"; + } + return "unknown"; +} + +struct options { + std::string hash_name = "mulshrolate3rx"; + std::string compile_mode_name = "ptx"; + std::string lookup_mode_name = "direct"; + std::string table_load_mode_name = "generic"; + std::string cpu_backend_name = "none"; + std::string keys_file; + std::string probe_keys_file; + std::string source_out_path; + int device_ordinal = 0; + int items_per_thread = 4; + int threads = 128; + int cpu_vector_width = 16; + int iterations = 10; + int warmup = 2; + std::uint64_t max_keys = 0; + std::uint64_t max_probe_keys = 0; + bool dump_fragment = false; + bool csv = false; + bool csv_header = false; + bool cpu_strict_vector_width = false; + bool analyze_slot_reuse = false; + bool analysis_only = false; + bool embed_table_data = false; + bool verify = true; +}; + +struct benchmark_result { + std::string key_source; + std::string probe_source; + std::string hash_name; + std::string device_name; + std::string cpu_backend_requested; + std::string cpu_backend_effective; + std::string cpu_key_mode; + compile_mode mode = compile_mode::ptx; + lookup_mode lookup = lookup_mode::direct; + table_load_mode table_load = table_load_mode::generic; + int device_ordinal = 0; + int items_per_thread = 0; + int threads = 0; + int blocks = 0; + int warmup = 0; + int iterations = 0; + int cpu_vector_requested = 0; + int cpu_vector_effective = 0; + int major = 0; + int minor = 0; + bool embedded_table_data = false; + bool cpu_enabled = false; + bool cpu_strict_vector_width = false; + bool cpu_jit_enabled = false; + bool slot_reuse_analyzed = false; + std::size_t key_count = 0; + std::size_t key_bytes = 0; + std::size_t probe_key_count = 0; + std::size_t probe_key_bytes = 0; + std::size_t table_data_bytes = 0; + std::size_t table_data_elements = 0; + std::size_t fragment_bytes = 0; + std::size_t combined_bytes = 0; + std::size_t image_bytes = 0; + std::uint32_t table_data_element_size = 0; + std::int32_t cpu_compile_hr = 0; + std::size_t host_rss_before_build_bytes = 0; + std::size_t host_rss_after_build_bytes = 0; + std::size_t host_rss_delta_bytes = 0; + std::size_t host_peak_rss_bytes = 0; + std::size_t vram_explicit_bytes = 0; + std::size_t vram_total_bytes = 0; + double build_ms = 0.0; + double table_export_ms = 0.0; + double emit_ms = 0.0; + double compose_ms = 0.0; + double compile_ms = 0.0; + double link_ms = 0.0; + double module_load_ms = 0.0; + double alloc_ms = 0.0; + double h2d_ms = 0.0; + double table_h2d_ms = 0.0; + double cpu_compile_ms = 0.0; + double cpu_lookup_avg_ms = 0.0; + double cpu_lookup_min_ms = 0.0; + double cpu_lookup_max_ms = 0.0; + double cpu_lookup_ns_per_key = 0.0; + double slot_compute_avg_ms = 0.0; + double slot_gather_avg_ms = 0.0; + double warp_unique_requests_avg = 0.0; + double warp_duplicate_ratio = 0.0; + double block_unique_requests_avg = 0.0; + double block_duplicate_ratio = 0.0; + double kernel_avg_ms = 0.0; + double kernel_min_ms = 0.0; + double kernel_max_ms = 0.0; + double gpu_lookup_ns_per_key = 0.0; + double slot_compute_ns_per_key = 0.0; + double slot_gather_ns_per_key = 0.0; + double d2h_ms = 0.0; + double verify_ms = 0.0; + std::uint32_t index_min = 0; + std::uint32_t index_max = 0; + std::size_t index_span = 0; +}; + +void print_usage(char const* argv0) +{ + std::cout << "Usage: " << argv0 + << " [--keys-file ] [--probe-keys-file ] [--max-keys ] [--max-probe-keys ] [--hash ]\n" + " [--items-per-thread ] [--threads ] [--iterations ]\n" + " [--warmup ] [--device ] [--compile-mode ]\n" + " [--lookup-mode ]\n" + " [--table-load-mode ]\n" + " [--cpu-backend ]\n" + " [--cpu-vector-width <1|2|4|8|16>] [--cpu-strict-vector-width <0|1>]\n" + " [--analyze-slot-reuse] [--analysis-only]\n" + " [--dump-fragment] [--source-out ] [--embed-table-data]\n" + " [--csv] [--csv-header]\n" + " [--no-verify]\n"; + std::cout << "Hash names: multiplyshiftr, multiplyshiftrx, mulshrolate1rx,\n"; + std::cout << " mulshrolate2rx, mulshrolate3rx, mulshrolate4rx\n"; +} + +double elapsed_ms(steady_clock::time_point start, steady_clock::time_point end) +{ + return std::chrono::duration(end - start).count(); +} + +std::string format_bytes(std::size_t bytes) +{ + static constexpr char const* suffixes[] = {"B", "KiB", "MiB", "GiB", "TiB"}; + double value = static_cast(bytes); + std::size_t suffix_index = 0; + while (value >= 1024.0 && suffix_index < (std::size(suffixes) - 1)) { + value /= 1024.0; + ++suffix_index; + } + + std::ostringstream stream; + stream << std::fixed << std::setprecision(suffix_index == 0 ? 0 : 2) + << value << suffixes[suffix_index]; + return stream.str(); +} + +std::string csv_escape(std::string const& value) +{ + std::string escaped = "\""; + escaped.reserve(value.size() + 2); + for (char c : value) { + if (c == '"') { escaped += "\"\""; } + else { escaped.push_back(c); } + } + escaped.push_back('"'); + return escaped; +} + +std::string cu_result_to_string(CUresult result) +{ + char const* name = nullptr; + char const* desc = nullptr; + cuGetErrorName(result, &name); + cuGetErrorString(result, &desc); + std::ostringstream stream; + stream << (name ? name : "CUDA_ERROR") << ": " << (desc ? desc : "unknown"); + return stream.str(); +} + +std::string nvrtc_result_to_string(nvrtcResult result) +{ + return nvrtcGetErrorString(result); +} + +std::string nvjitlink_result_to_string(nvJitLinkResult result) +{ + std::ostringstream stream; + stream << "nvJitLink(" << static_cast(result) << ")"; + return stream.str(); +} + +void throw_if_bad(CUresult result, char const* what) +{ + if (result != CUDA_SUCCESS) { + throw std::runtime_error(std::string(what) + " failed: " + + cu_result_to_string(result)); + } +} + +void throw_if_bad(cudaError_t result, char const* what) +{ + if (result != cudaSuccess) { + throw std::runtime_error(std::string(what) + " failed: " + + cudaGetErrorString(result)); + } +} + +void throw_if_bad(std::int32_t hr, char const* what) +{ + if (!succeeded(hr)) { + throw std::runtime_error(std::string(what) + " failed: " + + to_hex(static_cast(hr))); + } +} + +std::string get_nvrtc_log(nvrtcProgram program) +{ + std::size_t log_size = 0; + if (nvrtcGetProgramLogSize(program, &log_size) != NVRTC_SUCCESS || log_size == 0) { + return {}; + } + std::string log(log_size, '\0'); + if (nvrtcGetProgramLog(program, log.data()) != NVRTC_SUCCESS) { return {}; } + return log; +} + +std::string get_nvjitlink_log(nvJitLinkHandle handle, bool error_log) +{ + std::size_t log_size = 0; + nvJitLinkResult result = + error_log ? nvJitLinkGetErrorLogSize(handle, &log_size) + : nvJitLinkGetInfoLogSize(handle, &log_size); + if (result != NVJITLINK_SUCCESS || log_size == 0) { return {}; } + std::string log(log_size, '\0'); + result = error_log ? nvJitLinkGetErrorLog(handle, log.data()) + : nvJitLinkGetInfoLog(handle, log.data()); + if (result != NVJITLINK_SUCCESS) { return {}; } + return log; +} + +std::vector make_sample_keys() +{ + std::vector keys; + keys.reserve(64); + constexpr std::uint64_t high_bits = (1ULL << 40); + for (std::uint64_t key : {1ull, 3ull, 5ull, 7ull, 11ull, 13ull, 17ull, 19ull, + 23ull, 29ull, 31ull, 37ull, 41ull, 43ull, 47ull, 53ull, + 59ull, 61ull, 67ull, 71ull, 73ull, 79ull, 83ull, 89ull, + 97ull, 101ull, 103ull, 107ull, 109ull, 113ull, 127ull, 131ull, + 137ull, 139ull, 149ull, 151ull, 157ull, 163ull, 167ull, 173ull, + 179ull, 181ull, 191ull, 193ull, 197ull, 199ull, 211ull, 223ull, + 227ull, 229ull, 233ull, 239ull, 241ull, 251ull, 257ull, 263ull, + 269ull, 271ull, 277ull, 281ull, 283ull, 293ull, 307ull, 311ull}) { + keys.push_back(high_bits | key); + } + return keys; +} + +std::vector load_keys_file(std::string const& path, std::uint64_t max_keys) +{ + std::ifstream input(path, std::ios::binary | std::ios::ate); + if (!input) { + throw std::runtime_error("Unable to open keys file: " + path); + } + + auto const size = input.tellg(); + if (size < 0) { + throw std::runtime_error("Unable to determine keys file size: " + path); + } + if ((static_cast(size) % sizeof(std::uint64_t)) != 0) { + throw std::runtime_error("Keys file size is not a multiple of 8 bytes: " + path); + } + + auto const available_keys = + static_cast(size) / static_cast(sizeof(std::uint64_t)); + auto const keys_to_read = + (max_keys != 0 && max_keys < available_keys) ? max_keys : available_keys; + + std::vector keys(static_cast(keys_to_read)); + input.seekg(0, std::ios::beg); + input.read(reinterpret_cast(keys.data()), + static_cast(keys.size() * sizeof(std::uint64_t))); + if (!input) { + throw std::runtime_error("Unable to read keys file contents: " + path); + } + + return keys; +} + +double ns_per_key(double milliseconds, std::size_t key_count) +{ + if (key_count == 0) { return 0.0; } + return (milliseconds * 1'000'000.0) / static_cast(key_count); +} + +std::size_t current_rss_bytes() +{ +#if defined(__linux__) + long page_size = sysconf(_SC_PAGESIZE); + std::ifstream input("/proc/self/statm"); + std::size_t total_pages = 0; + std::size_t resident_pages = 0; + if (input >> total_pages >> resident_pages) { + return resident_pages * static_cast(page_size); + } +#endif + return 0; +} + +std::size_t peak_rss_bytes() +{ +#if !defined(_WIN32) + struct rusage usage {}; + if (getrusage(RUSAGE_SELF, &usage) == 0) { +#if defined(__APPLE__) + return static_cast(usage.ru_maxrss); +#else + return static_cast(usage.ru_maxrss) * 1024ull; +#endif + } +#endif + return 0; +} + +std::uint64_t extract_bits64(std::uint64_t value, std::uint64_t bitmap) +{ + std::uint64_t result = 0; + std::uint64_t out_bit = 0; + while (bitmap != 0) { + auto const lsb = bitmap & (~bitmap + 1); + if ((value & lsb) != 0) { + result |= (1ULL << out_bit); + } + bitmap ^= lsb; + ++out_bit; + } + return result; +} + +std::uint64_t mask_from_bits(std::uint32_t bits) +{ + return (bits >= 64u) ? std::numeric_limits::max() + : ((1ULL << bits) - 1ULL); +} + +std::uint32_t downsize_key64(std::uint64_t key, PH_ONLINE_JIT_TABLE_INFO const& info) +{ + auto const original_key_mask = mask_from_bits(info.OriginalKeySizeInBytes * 8u); + auto const key_mask = mask_from_bits(info.KeySizeInBytes * 8u); + key &= original_key_mask; + if (info.DownsizeBitmap != 0) { + return static_cast(extract_bits64(key, info.DownsizeBitmap) & key_mask); + } + return static_cast(key & key_mask); +} + +std::vector make_downsized_keys32(std::vector const& keys, + PH_ONLINE_JIT_TABLE_INFO const& info) +{ + std::vector downsized; + downsized.reserve(keys.size()); + for (auto key : keys) { + downsized.push_back(downsize_key64(key, info)); + } + return downsized; +} + +std::uint32_t rotr32_host(std::uint32_t value, std::uint32_t shift) +{ + shift &= 31u; + if (shift == 0u) { return value; } + return (value >> shift) | (value << (32u - shift)); +} + +struct host_slot_pair { + std::uint32_t first = 0; + std::uint32_t second = 0; +}; + +std::uint32_t table_value_from_host_data(void const* table_data, + PH_ONLINE_JIT_TABLE_INFO const& info, + std::uint32_t index) +{ + if (info.AssignedElementSizeInBytes == sizeof(std::uint16_t)) { + auto const* values = static_cast(table_data); + return static_cast(values[index]); + } + auto const* values = static_cast(table_data); + return values[index]; +} + +host_slot_pair slot_pair_from_key_host(std::uint64_t key, + PH_ONLINE_JIT_HASH_FUNCTION hash_function, + PH_ONLINE_JIT_TABLE_INFO const& info) +{ + auto const downsized = downsize_key64(key, info); + auto const seed1 = static_cast(info.Seed1); + auto const seed2 = static_cast(info.Seed2); + auto const seed4 = static_cast(info.Seed4); + auto const seed5 = static_cast(info.Seed5); + auto const seed3_byte1 = (info.Seed3 & 0xffu); + auto const seed3_byte2 = ((info.Seed3 >> 8u) & 0xffu); + auto const seed3_byte3 = ((info.Seed3 >> 16u) & 0xffu); + auto const seed3_byte4 = ((info.Seed3 >> 24u) & 0xffu); + auto const use_32bit_math = (info.KeySizeInBytes <= sizeof(std::uint32_t)); + + if (use_32bit_math) { + auto const downsized32 = static_cast(downsized); + auto const seed1_32 = static_cast(seed1); + auto const seed2_32 = static_cast(seed2); + auto const seed4_32 = static_cast(seed4); + auto const seed5_32 = static_cast(seed5); + + switch (hash_function) { + case PhOnlineJitHashMultiplyShiftR: { + auto const vertex1 = (downsized32 * seed1_32) >> seed3_byte1; + auto const vertex2 = (downsized32 * seed2_32) >> seed3_byte2; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftLR: { + auto const vertex1 = (downsized32 * seed1_32) << seed3_byte1; + auto const vertex2 = (downsized32 * seed2_32) >> seed3_byte2; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftRMultiply: { + auto const vertex1 = ((downsized32 * seed1_32) >> seed3_byte1) * seed2_32; + auto const vertex2 = ((downsized32 * seed4_32) >> seed3_byte2) * seed5_32; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftR2: { + auto const vertex1 = (((downsized32 * seed1_32) >> seed3_byte1) * seed2_32) >> seed3_byte2; + auto const vertex2 = (((downsized32 * seed4_32) >> seed3_byte3) * seed5_32) >> seed3_byte4; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftRX: { + auto const vertex1 = (downsized32 * seed1_32) >> seed3_byte1; + auto const vertex2 = (downsized32 * seed2_32) >> seed3_byte1; + return {vertex1, vertex2}; + } + case PhOnlineJitHashMulshrolate1RX: { + auto vertex1 = rotr32_host(downsized32 * seed1_32, seed3_byte2); + vertex1 >>= seed3_byte1; + auto vertex2 = downsized32 * seed2_32; + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + case PhOnlineJitHashMulshrolate2RX: { + auto vertex1 = rotr32_host(downsized32 * seed1_32, seed3_byte2); + vertex1 >>= seed3_byte1; + auto vertex2 = rotr32_host(downsized32 * seed2_32, seed3_byte3); + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + case PhOnlineJitHashMulshrolate3RX: { + auto vertex1 = rotr32_host(downsized32 * seed1_32, seed3_byte2); + vertex1 = vertex1 * seed4_32; + vertex1 >>= seed3_byte1; + auto vertex2 = rotr32_host(downsized32 * seed2_32, seed3_byte3); + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + case PhOnlineJitHashMulshrolate4RX: { + auto vertex1 = rotr32_host(downsized32 * seed1_32, seed3_byte2); + vertex1 = vertex1 * seed4_32; + vertex1 >>= seed3_byte1; + auto vertex2 = rotr32_host(downsized32 * seed2_32, seed3_byte3); + vertex2 = vertex2 * seed5_32; + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + } + } + + auto const downsized64 = static_cast(downsized); + + switch (hash_function) { + case PhOnlineJitHashMultiplyShiftR: { + auto vertex1 = (downsized64 * seed1) >> seed3_byte1; + auto vertex2 = (downsized64 * seed2) >> seed3_byte2; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftLR: { + auto vertex1 = (downsized64 * seed1) << seed3_byte1; + auto vertex2 = (downsized64 * seed2) >> seed3_byte2; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftRMultiply: { + auto vertex1 = ((downsized64 * seed1) >> seed3_byte1) * seed2; + auto vertex2 = ((downsized64 * seed4) >> seed3_byte2) * seed5; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftR2: { + auto vertex1 = (((downsized64 * seed1) >> seed3_byte1) * seed2) >> seed3_byte2; + auto vertex2 = (((downsized64 * seed4) >> seed3_byte3) * seed5) >> seed3_byte4; + return {static_cast(vertex1 & info.HashMask), + static_cast(vertex2 & info.HashMask)}; + } + case PhOnlineJitHashMultiplyShiftRX: { + auto vertex1 = (downsized64 * seed1) >> seed3_byte1; + auto vertex2 = (downsized64 * seed2) >> seed3_byte1; + return {static_cast(vertex1), + static_cast(vertex2)}; + } + case PhOnlineJitHashMulshrolate1RX: { + auto downsized32 = static_cast(downsized64); + auto vertex1 = rotr32_host(downsized32 * static_cast(seed1), seed3_byte2); + vertex1 >>= seed3_byte1; + auto vertex2 = downsized32 * static_cast(seed2); + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + case PhOnlineJitHashMulshrolate2RX: { + auto downsized32 = static_cast(downsized); + auto vertex1 = rotr32_host(downsized32 * static_cast(seed1), seed3_byte2); + vertex1 >>= seed3_byte1; + auto vertex2 = rotr32_host(downsized32 * static_cast(seed2), seed3_byte3); + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + case PhOnlineJitHashMulshrolate3RX: { + auto downsized32 = static_cast(downsized); + auto vertex1 = rotr32_host(downsized32 * static_cast(seed1), seed3_byte2); + vertex1 = vertex1 * static_cast(seed4); + vertex1 >>= seed3_byte1; + auto vertex2 = rotr32_host(downsized32 * static_cast(seed2), seed3_byte3); + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + case PhOnlineJitHashMulshrolate4RX: { + auto downsized32 = static_cast(downsized); + auto vertex1 = rotr32_host(downsized32 * static_cast(seed1), seed3_byte2); + vertex1 = vertex1 * static_cast(seed4); + vertex1 >>= seed3_byte1; + auto vertex2 = rotr32_host(downsized32 * static_cast(seed2), seed3_byte3); + vertex2 = vertex2 * static_cast(seed5); + vertex2 >>= seed3_byte1; + return {vertex1, vertex2}; + } + } + + throw std::runtime_error("Unsupported hash function for host slot analysis"); +} + +std::uint32_t index_from_key_host(std::uint64_t key, + PH_ONLINE_JIT_HASH_FUNCTION hash_function, + PH_ONLINE_JIT_TABLE_INFO const& info, + void const* table_data) +{ + auto const slots = slot_pair_from_key_host(key, hash_function, info); + if (slots.first >= info.NumberOfTableElements || slots.second >= info.NumberOfTableElements) { + throw std::runtime_error("Host slot pair exceeded exported table bounds"); + } + auto const low = table_value_from_host_data(table_data, info, slots.first); + auto const high = table_value_from_host_data(table_data, info, slots.second); + return static_cast((low + high) & info.IndexMask); +} + +struct tile_reuse_stats { + double unique_requests_avg = 0.0; + double duplicate_ratio = 0.0; +}; + +tile_reuse_stats analyze_slot_reuse_for_tile(std::vector const& keys, + PH_ONLINE_JIT_HASH_FUNCTION hash_function, + PH_ONLINE_JIT_TABLE_INFO const& info, + std::size_t threads_per_tile, + std::size_t items_per_thread) +{ + tile_reuse_stats stats; + if (keys.empty() || threads_per_tile == 0 || items_per_thread == 0) { + return stats; + } + + auto const keys_per_tile = threads_per_tile * items_per_thread; + std::vector requests; + requests.reserve(keys_per_tile * 2u); + + std::size_t total_tiles = 0; + std::size_t total_requests = 0; + std::size_t total_unique = 0; + + for (std::size_t base = 0; base < keys.size(); base += keys_per_tile) { + requests.clear(); + auto const limit = std::min(base + keys_per_tile, keys.size()); + for (std::size_t i = base; i < limit; ++i) { + auto const slots = slot_pair_from_key_host(keys[i], hash_function, info); + requests.push_back(slots.first); + requests.push_back(slots.second); + } + std::sort(requests.begin(), requests.end()); + auto const unique_end = std::unique(requests.begin(), requests.end()); + auto const unique_count = static_cast(unique_end - requests.begin()); + total_tiles += 1; + total_requests += requests.size(); + total_unique += unique_count; + } + + if (total_tiles != 0) { + stats.unique_requests_avg = static_cast(total_unique) / + static_cast(total_tiles); + } + if (total_requests != 0) { + stats.duplicate_ratio = 1.0 - + (static_cast(total_unique) / static_cast(total_requests)); + } + + return stats; +} + +std::string compose_translation_unit(std::string const& fragment, + int items_per_thread, + int threads_per_block, + bool embed_table_data, + bool validate_membership, + lookup_mode mode, + table_load_mode table_load) +{ + if (validate_membership && mode != lookup_mode::direct) { + throw std::runtime_error( + "compose_translation_unit only supports membership validation for lookup-mode=direct"); + } + + std::ostringstream source; + source << fragment << "\n"; + source << R"( +namespace perfecthash::consumer { +namespace generated = perfecthash::generated::online_jit_table; + +__device__ __forceinline__ uint32_t load_table_value( + const generated::table_data_type* table_data, + uint32_t index) +{ +)"; + if (table_load == table_load_mode::readonly) { + source << R"(#if defined(__CUDA_ARCH__) + return static_cast(__ldg(table_data + index)); +#else + return static_cast(table_data[index]); +#endif +)"; + } else { + source << R"( return static_cast(table_data[index]); +)"; + } + source << R"(} + +template +__device__ __forceinline__ void direct_probe_tile( + const generated::original_key_type (&input)[ITEMS_PER_THREAD], + const bool (&valid)[ITEMS_PER_THREAD], + const generated::table_data_type* table_data, + const generated::original_key_type* build_keys, + size_t build_key_count, + uint32_t (&output)[ITEMS_PER_THREAD]) +{ +#pragma unroll + for (int item = 0; item < ITEMS_PER_THREAD; ++item) { + if (!valid[item]) { + output[item] = 0; + continue; + } + const generated::slot_pair_type slots = generated::slot_pair_from_key(input[item]); + if (slots.first >= generated::number_of_table_elements || + slots.second >= generated::number_of_table_elements) { + output[item] = 0xFFFFFFFFu; + continue; + } + const uint32_t value_low = load_table_value(table_data, slots.first); + const uint32_t value_high = load_table_value(table_data, slots.second); + const uint32_t index = static_cast((value_low + value_high) & generated::index_mask); +)"; + if (validate_membership) { + source << R"( output[item] = + (index < build_key_count && build_keys[index] == input[item]) ? index : 0xFFFFFFFFu; +)"; + } else { + source << R"( output[item] = index; +)"; + } + source << R"( } +} + +__device__ __forceinline__ uint32_t warp_cached_load( + const generated::table_data_type* table_data, + uint32_t slot) +{ + const unsigned active = __activemask(); + const unsigned match = __match_any_sync(active, slot); + const int lane = static_cast(threadIdx.x & 31); + const int leader = __ffs(static_cast(match)) - 1; + uint32_t value = 0; + if (lane == leader) { + value = load_table_value(table_data, slot); + } + return __shfl_sync(active, value, leader); +} + +template +__device__ __forceinline__ void warp_cached_probe_tile( + const generated::original_key_type (&input)[ITEMS_PER_THREAD], + const bool (&valid)[ITEMS_PER_THREAD], + const generated::table_data_type* table_data, + uint32_t (&output)[ITEMS_PER_THREAD]) +{ +#pragma unroll + for (int item = 0; item < ITEMS_PER_THREAD; ++item) { + if (!valid[item]) { + output[item] = 0; + continue; + } + const generated::slot_pair_type slots = generated::slot_pair_from_key(input[item]); + if (slots.first >= generated::number_of_table_elements || + slots.second >= generated::number_of_table_elements) { + output[item] = 0xFFFFFFFFu; + continue; + } + const uint32_t value_low = warp_cached_load(table_data, slots.first); + const uint32_t value_high = warp_cached_load(table_data, slots.second); + output[item] = static_cast((value_low + value_high) & generated::index_mask); + } +} + +template +__device__ __forceinline__ void compute_slot_tile( + const generated::original_key_type (&input)[ITEMS_PER_THREAD], + const bool (&valid)[ITEMS_PER_THREAD], + uint32_t (&slot1)[ITEMS_PER_THREAD], + uint32_t (&slot2)[ITEMS_PER_THREAD]) +{ +#pragma unroll + for (int item = 0; item < ITEMS_PER_THREAD; ++item) { + if (!valid[item]) { + slot1[item] = 0; + slot2[item] = 0; + continue; + } + const generated::slot_pair_type slots = generated::slot_pair_from_key(input[item]); + slot1[item] = slots.first; + slot2[item] = slots.second; + } +} + +template +__device__ __forceinline__ void gather_tile( + const uint32_t (&slot1)[ITEMS_PER_THREAD], + const uint32_t (&slot2)[ITEMS_PER_THREAD], + const bool (&valid)[ITEMS_PER_THREAD], + const generated::table_data_type* table_data, + uint32_t (&output)[ITEMS_PER_THREAD]) +{ +#pragma unroll + for (int item = 0; item < ITEMS_PER_THREAD; ++item) { + if (!valid[item]) { + output[item] = 0; + continue; + } + if (slot1[item] >= generated::number_of_table_elements || + slot2[item] >= generated::number_of_table_elements) { + output[item] = 0xFFFFFFFFu; + continue; + } + const uint32_t value_low = load_table_value(table_data, slot1[item]); + const uint32_t value_high = load_table_value(table_data, slot2[item]); + output[item] = static_cast((value_low + value_high) & generated::index_mask); + } +} + +__device__ __forceinline__ void swap_request( + uint32_t& slot_a, + uint32_t& slot_b, + uint32_t& dest_a, + uint32_t& dest_b) +{ + const uint32_t slot_tmp = slot_a; + slot_a = slot_b; + slot_b = slot_tmp; + const uint32_t dest_tmp = dest_a; + dest_a = dest_b; + dest_b = dest_tmp; +} + +} // namespace perfecthash::consumer +)"; + if (mode == lookup_mode::direct) { + source << R"( +extern "C" __global__ void probe_kernel( + const perfecthash::generated::online_jit_table::original_key_type* keys, +)"; + if (validate_membership) { + source << R"( const perfecthash::generated::online_jit_table::original_key_type* build_keys, + size_t build_key_count, +)"; + } + if (!embed_table_data) { + source << R"( const perfecthash::generated::online_jit_table::table_data_type* table_data, +)"; + } + source << R"( uint32_t* out, + size_t count) +{ + namespace generated = perfecthash::generated::online_jit_table; + constexpr int ITEMS = )"; + source << items_per_thread; + source << R"(; + const size_t thread_base = + (static_cast(blockIdx.x) * blockDim.x + threadIdx.x) * ITEMS; + + generated::original_key_type local_keys[ITEMS] = {}; + bool local_valid[ITEMS] = {}; + uint32_t local_out[ITEMS] = {}; + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { + local_keys[item] = keys[idx]; + local_valid[item] = true; + } + } + + perfecthash::consumer::direct_probe_tile(local_keys, local_valid, )"; + if (embed_table_data) { + source << "generated::table_data"; + } else { + source << "table_data"; + } + source << ", "; + if (validate_membership) { + source << "build_keys, build_key_count"; + } else { + source << "nullptr, 0"; + } + source << R"(, local_out); + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { out[idx] = local_out[item]; } + } +} +)"; + } else if (mode == lookup_mode::split) { + source << R"( +extern "C" __global__ void compute_slots_kernel( + const perfecthash::generated::online_jit_table::original_key_type* keys, + uint32_t* slot1_out, + uint32_t* slot2_out, + size_t count) +{ + namespace generated = perfecthash::generated::online_jit_table; + constexpr int ITEMS = )"; + source << items_per_thread; + source << R"(; + const size_t thread_base = + (static_cast(blockIdx.x) * blockDim.x + threadIdx.x) * ITEMS; + + generated::original_key_type local_keys[ITEMS] = {}; + bool local_valid[ITEMS] = {}; + uint32_t local_slot1[ITEMS] = {}; + uint32_t local_slot2[ITEMS] = {}; + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { + local_keys[item] = keys[idx]; + local_valid[item] = true; + } + } + + perfecthash::consumer::compute_slot_tile(local_keys, local_valid, local_slot1, local_slot2); + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { + slot1_out[idx] = local_slot1[item]; + slot2_out[idx] = local_slot2[item]; + } + } +} + +extern "C" __global__ void gather_kernel( + const uint32_t* slot1_in, + const uint32_t* slot2_in, +)"; + if (!embed_table_data) { + source << R"( const perfecthash::generated::online_jit_table::table_data_type* table_data, +)"; + } + source << R"( uint32_t* out, + size_t count) +{ + namespace generated = perfecthash::generated::online_jit_table; + constexpr int ITEMS = )"; + source << items_per_thread; + source << R"(; + const size_t thread_base = + (static_cast(blockIdx.x) * blockDim.x + threadIdx.x) * ITEMS; + + uint32_t local_slot1[ITEMS] = {}; + uint32_t local_slot2[ITEMS] = {}; + bool local_valid[ITEMS] = {}; + uint32_t local_out[ITEMS] = {}; + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { + local_slot1[item] = slot1_in[idx]; + local_slot2[item] = slot2_in[idx]; + local_valid[item] = true; + } + } + + perfecthash::consumer::gather_tile(local_slot1, local_slot2, local_valid, )"; + if (embed_table_data) { + source << "generated::table_data"; + } else { + source << "table_data"; + } + source << R"(, local_out); + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { out[idx] = local_out[item]; } + } +} +)"; + } else if (mode == lookup_mode::warpcache) { + source << R"( +extern "C" __global__ void warpcache_probe_kernel( + const perfecthash::generated::online_jit_table::original_key_type* keys, +)"; + if (!embed_table_data) { + source << R"( const perfecthash::generated::online_jit_table::table_data_type* table_data, +)"; + } + source << R"( uint32_t* out, + size_t count) +{ + namespace generated = perfecthash::generated::online_jit_table; + constexpr int ITEMS = )"; + source << items_per_thread; + source << R"(; + const size_t thread_base = + (static_cast(blockIdx.x) * blockDim.x + threadIdx.x) * ITEMS; + + generated::original_key_type local_keys[ITEMS] = {}; + bool local_valid[ITEMS] = {}; + uint32_t local_out[ITEMS] = {}; + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { + local_keys[item] = keys[idx]; + local_valid[item] = true; + } + } + + perfecthash::consumer::warp_cached_probe_tile(local_keys, local_valid, )"; + if (embed_table_data) { + source << "generated::table_data"; + } else { + source << "table_data"; + } + source << R"(, local_out); + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t idx = thread_base + static_cast(item); + if (idx < count) { out[idx] = local_out[item]; } + } +} +)"; + } else { + source << R"( +extern "C" __global__ void blocksort_probe_kernel( + const perfecthash::generated::online_jit_table::original_key_type* keys, +)"; + if (!embed_table_data) { + source << R"( const perfecthash::generated::online_jit_table::table_data_type* table_data, +)"; + } + source << R"( uint32_t* out, + size_t count) +{ + namespace generated = perfecthash::generated::online_jit_table; + constexpr int ITEMS = )"; + source << items_per_thread; + source << R"(; + constexpr int THREADS = )"; + source << threads_per_block; + source << R"(; + constexpr int TOTAL_ITEMS = THREADS * ITEMS; + constexpr int TOTAL_REQUESTS = TOTAL_ITEMS * 2; + constexpr uint32_t INVALID_SLOT = 0xFFFFFFFFu; + + __shared__ uint32_t shared_slots[TOTAL_REQUESTS]; + __shared__ uint32_t shared_dests[TOTAL_REQUESTS]; + __shared__ uint32_t shared_values[TOTAL_REQUESTS]; + __shared__ uint32_t shared_partials[TOTAL_REQUESTS]; + + const uint32_t tid = threadIdx.x; + const size_t block_base = static_cast(blockIdx.x) * static_cast(TOTAL_ITEMS); + const uint32_t local_base = tid * ITEMS; + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const size_t global_idx = block_base + static_cast(local_base + item); + const uint32_t request_base = (local_base + item) * 2u; + if (global_idx < count) { + const generated::slot_pair_type slots = generated::slot_pair_from_key(keys[global_idx]); + shared_slots[request_base] = slots.first; + shared_slots[request_base + 1] = slots.second; + shared_dests[request_base] = request_base; + shared_dests[request_base + 1] = request_base + 1u; + shared_values[request_base] = 0; + shared_values[request_base + 1] = 0; + shared_partials[request_base] = 0; + shared_partials[request_base + 1] = 0; + } else { + shared_slots[request_base] = INVALID_SLOT; + shared_slots[request_base + 1] = INVALID_SLOT; + shared_dests[request_base] = request_base; + shared_dests[request_base + 1] = request_base + 1u; + shared_values[request_base] = 0; + shared_values[request_base + 1] = 0; + shared_partials[request_base] = 0; + shared_partials[request_base + 1] = 0; + } + } + __syncthreads(); + + for (uint32_t k = 2; k <= TOTAL_REQUESTS; k <<= 1) { + for (uint32_t j = k >> 1; j > 0; j >>= 1) { + for (uint32_t i = tid; i < TOTAL_REQUESTS; i += THREADS) { + const uint32_t ixj = i ^ j; + if (ixj > i) { + const bool i_invalid = (shared_slots[i] == INVALID_SLOT); + const bool ixj_invalid = (shared_slots[ixj] == INVALID_SLOT); + const bool ascending = ((i & k) == 0); + bool should_swap = false; + if (i_invalid != ixj_invalid) { + should_swap = i_invalid && !ixj_invalid; + } else if (!i_invalid) { + should_swap = + ascending ? (shared_slots[i] > shared_slots[ixj]) + : (shared_slots[i] < shared_slots[ixj]); + } + if (should_swap) { + perfecthash::consumer::swap_request(shared_slots[i], + shared_slots[ixj], + shared_dests[i], + shared_dests[ixj]); + } + } + } + __syncthreads(); + } + } + + for (uint32_t i = tid; i < TOTAL_REQUESTS; i += THREADS) { + const uint32_t slot = shared_slots[i]; + if (slot != INVALID_SLOT) { +)"; + if (embed_table_data) { + source << R"( if (i == 0 || shared_slots[i - 1] != slot) { + shared_values[i] = + perfecthash::consumer::load_table_value(generated::table_data, slot); + } +)"; + } else { + source << R"( if (i == 0 || shared_slots[i - 1] != slot) { + shared_values[i] = + perfecthash::consumer::load_table_value(table_data, slot); + } +)"; + } + source << R"( + } + } + __syncthreads(); + + for (uint32_t i = tid; i < TOTAL_REQUESTS; i += THREADS) { + const uint32_t slot = shared_slots[i]; + if (slot != INVALID_SLOT && i > 0 && shared_slots[i - 1] == slot) { + uint32_t first = i - 1; + while (first > 0 && shared_slots[first - 1] == slot) { + --first; + } + shared_values[i] = shared_values[first]; + } + } + __syncthreads(); + + for (uint32_t i = tid; i < TOTAL_REQUESTS; i += THREADS) { + const uint32_t slot = shared_slots[i]; + if (slot != INVALID_SLOT) { + shared_partials[shared_dests[i]] = shared_values[i]; + } + } + __syncthreads(); + +#pragma unroll + for (int item = 0; item < ITEMS; ++item) { + const uint32_t output_base = (local_base + item) * 2u; + const size_t global_idx = block_base + static_cast(local_base + item); + if (global_idx < count) { + const uint32_t value_low = shared_partials[output_base]; + const uint32_t value_high = shared_partials[output_base + 1u]; + out[global_idx] = static_cast((value_low + value_high) & generated::index_mask); + } + } +} +)"; + } + return source.str(); +} + +std::string compile_to_ptx(std::string const& source, int major, int minor) +{ + nvrtcProgram program{}; + auto create_result = + nvrtcCreateProgram(&program, source.c_str(), "perfecthash_online_nvrtc.cu", 0, nullptr, nullptr); + if (create_result != NVRTC_SUCCESS) { + throw std::runtime_error("nvrtcCreateProgram failed: " + + nvrtc_result_to_string(create_result)); + } + + std::string arch = "--gpu-architecture=compute_" + std::to_string(major) + std::to_string(minor); + std::string include_dir = std::string("--include-path=") + PH_CUDA_INCLUDE_DIR; + + std::vector options{ + "--std=c++17", + "--device-as-default-execution-space", + arch.c_str(), + include_dir.c_str(), + }; + + auto compile_result = + nvrtcCompileProgram(program, static_cast(options.size()), options.data()); + auto const log = get_nvrtc_log(program); + if (!log.empty()) { std::cerr << log; } + if (compile_result != NVRTC_SUCCESS) { + nvrtcDestroyProgram(&program); + throw std::runtime_error("nvrtcCompileProgram failed: " + + nvrtc_result_to_string(compile_result)); + } + + std::size_t ptx_size = 0; + if (nvrtcGetPTXSize(program, &ptx_size) != NVRTC_SUCCESS) { + nvrtcDestroyProgram(&program); + throw std::runtime_error("nvrtcGetPTXSize failed"); + } + + std::string ptx(ptx_size, '\0'); + if (nvrtcGetPTX(program, ptx.data()) != NVRTC_SUCCESS) { + nvrtcDestroyProgram(&program); + throw std::runtime_error("nvrtcGetPTX failed"); + } + + nvrtcDestroyProgram(&program); + return ptx; +} + +std::vector compile_to_ltoir(std::string const& source, int major, int minor) +{ + nvrtcProgram program{}; + auto create_result = + nvrtcCreateProgram(&program, source.c_str(), "perfecthash_online_nvrtc_lto.cu", 0, nullptr, nullptr); + if (create_result != NVRTC_SUCCESS) { + throw std::runtime_error("nvrtcCreateProgram failed: " + + nvrtc_result_to_string(create_result)); + } + + std::string arch = "--gpu-architecture=compute_" + std::to_string(major) + std::to_string(minor); + std::string include_dir = std::string("--include-path=") + PH_CUDA_INCLUDE_DIR; + + std::vector options{ + "--std=c++17", + "--device-as-default-execution-space", + "--dlink-time-opt", + "--gen-opt-lto", + arch.c_str(), + include_dir.c_str(), + }; + + auto compile_result = + nvrtcCompileProgram(program, static_cast(options.size()), options.data()); + auto const log = get_nvrtc_log(program); + if (!log.empty()) { std::cerr << log; } + if (compile_result != NVRTC_SUCCESS) { + nvrtcDestroyProgram(&program); + throw std::runtime_error("nvrtcCompileProgram failed: " + + nvrtc_result_to_string(compile_result)); + } + + std::size_t ir_size = 0; + if (nvrtcGetLTOIRSize(program, &ir_size) != NVRTC_SUCCESS) { + nvrtcDestroyProgram(&program); + throw std::runtime_error("nvrtcGetLTOIRSize failed"); + } + + std::vector ir(ir_size); + if (nvrtcGetLTOIR(program, ir.data()) != NVRTC_SUCCESS) { + nvrtcDestroyProgram(&program); + throw std::runtime_error("nvrtcGetLTOIR failed"); + } + + nvrtcDestroyProgram(&program); + return ir; +} + +std::vector link_ltoir_to_cubin(std::vector const& ir, int major, int minor) +{ + nvJitLinkHandle handle{}; + std::string arch = "-arch=sm_" + std::to_string(major) + std::to_string(minor); + std::vector options{ + arch.c_str(), + "-lto", + }; + + auto result = nvJitLinkCreate(&handle, + static_cast(options.size()), + options.data()); + if (result != NVJITLINK_SUCCESS) { + throw std::runtime_error("nvJitLinkCreate failed: " + + nvjitlink_result_to_string(result)); + } + + result = nvJitLinkAddData(handle, + NVJITLINK_INPUT_LTOIR, + ir.data(), + ir.size(), + "perfecthash_online_nvrtc_lto"); + if (result != NVJITLINK_SUCCESS) { + auto const error_log = get_nvjitlink_log(handle, true); + nvJitLinkDestroy(&handle); + throw std::runtime_error("nvJitLinkAddData failed: " + + nvjitlink_result_to_string(result) + + (error_log.empty() ? std::string{} : (": " + error_log))); + } + + result = nvJitLinkComplete(handle); + auto const info_log = get_nvjitlink_log(handle, false); + auto const error_log = get_nvjitlink_log(handle, true); + if (!info_log.empty()) { std::cerr << info_log; } + if (!error_log.empty()) { std::cerr << error_log; } + if (result != NVJITLINK_SUCCESS) { + nvJitLinkDestroy(&handle); + throw std::runtime_error("nvJitLinkComplete failed: " + + nvjitlink_result_to_string(result)); + } + + std::size_t cubin_size = 0; + result = nvJitLinkGetLinkedCubinSize(handle, &cubin_size); + if (result != NVJITLINK_SUCCESS) { + nvJitLinkDestroy(&handle); + throw std::runtime_error("nvJitLinkGetLinkedCubinSize failed: " + + nvjitlink_result_to_string(result)); + } + + std::vector cubin(cubin_size); + result = nvJitLinkGetLinkedCubin(handle, cubin.data()); + nvJitLinkDestroy(&handle); + if (result != NVJITLINK_SUCCESS) { + throw std::runtime_error("nvJitLinkGetLinkedCubin failed: " + + nvjitlink_result_to_string(result)); + } + + return cubin; +} + +struct ph_context_handle { + PH_ONLINE_JIT_CONTEXT* value = nullptr; + ~ph_context_handle() + { + if (value != nullptr) { PhOnlineJitClose(value); } + } +}; + +struct ph_table_handle { + PH_ONLINE_JIT_TABLE* value = nullptr; + ~ph_table_handle() + { + if (value != nullptr) { PhOnlineJitReleaseTable(value); } + } +}; + +struct cuda_source_handle { + char* value = nullptr; + ~cuda_source_handle() + { + if (value != nullptr) { PhOnlineJitFreeCudaSource(value); } + } +}; + +struct table_data_handle { + void* value = nullptr; + ~table_data_handle() + { + if (value != nullptr) { PhOnlineJitFreeCudaTableData(value); } + } +}; + +struct primary_context_handle { + CUdevice device{}; + bool retained = false; + ~primary_context_handle() + { + if (retained) { cuDevicePrimaryCtxRelease(device); } + } +}; + +struct module_handle { + CUmodule value{}; + ~module_handle() + { + if (value != nullptr) { cuModuleUnload(value); } + } +}; + +struct device_memory { + CUdeviceptr value{}; + ~device_memory() + { + if (value != 0) { cuMemFree(value); } + } +}; + +struct event_handle { + CUevent value{}; + ~event_handle() + { + if (value != nullptr) { cuEventDestroy(value); } + } +}; + +std::vector run_probe_iterations(CUfunction function, + void** params, + int blocks, + int threads, + int warmup, + int iterations) +{ + event_handle start_event; + event_handle stop_event; + throw_if_bad(cuEventCreate(&start_event.value, CU_EVENT_DEFAULT), "cuEventCreate(start)"); + throw_if_bad(cuEventCreate(&stop_event.value, CU_EVENT_DEFAULT), "cuEventCreate(stop)"); + + for (int iteration = 0; iteration < warmup; ++iteration) { + throw_if_bad(cuLaunchKernel(function, + static_cast(blocks), + 1, + 1, + static_cast(threads), + 1, + 1, + 0, + 0, + params, + nullptr), + "cuLaunchKernel(warmup)"); + throw_if_bad(cuCtxSynchronize(), "cuCtxSynchronize(warmup)"); + } + + std::vector timings; + timings.reserve(static_cast(iterations)); + for (int iteration = 0; iteration < iterations; ++iteration) { + throw_if_bad(cuEventRecord(start_event.value, 0), "cuEventRecord(start)"); + throw_if_bad(cuLaunchKernel(function, + static_cast(blocks), + 1, + 1, + static_cast(threads), + 1, + 1, + 0, + 0, + params, + nullptr), + "cuLaunchKernel"); + throw_if_bad(cuEventRecord(stop_event.value, 0), "cuEventRecord(stop)"); + throw_if_bad(cuEventSynchronize(stop_event.value), "cuEventSynchronize(stop)"); + + float kernel_ms = 0.0f; + throw_if_bad(cuEventElapsedTime(&kernel_ms, start_event.value, stop_event.value), + "cuEventElapsedTime"); + timings.push_back(static_cast(kernel_ms)); + } + + return timings; +} + +struct timing_stats { + double avg = 0.0; + double min = 0.0; + double max = 0.0; +}; + +timing_stats summarize_timings(std::vector const& timings); + +struct split_timing_stats { + timing_stats total; + timing_stats compute; + timing_stats gather; +}; + +split_timing_stats run_split_probe_iterations(CUfunction compute_function, + void** compute_params, + CUfunction gather_function, + void** gather_params, + int blocks, + int threads, + int warmup, + int iterations) +{ + event_handle start_event; + event_handle compute_stop_event; + event_handle stop_event; + throw_if_bad(cuEventCreate(&start_event.value, CU_EVENT_DEFAULT), "cuEventCreate(start)"); + throw_if_bad(cuEventCreate(&compute_stop_event.value, CU_EVENT_DEFAULT), "cuEventCreate(compute stop)"); + throw_if_bad(cuEventCreate(&stop_event.value, CU_EVENT_DEFAULT), "cuEventCreate(stop)"); + + for (int iteration = 0; iteration < warmup; ++iteration) { + throw_if_bad(cuLaunchKernel(compute_function, + static_cast(blocks), + 1, + 1, + static_cast(threads), + 1, + 1, + 0, + 0, + compute_params, + nullptr), + "cuLaunchKernel(split warmup compute)"); + throw_if_bad(cuLaunchKernel(gather_function, + static_cast(blocks), + 1, + 1, + static_cast(threads), + 1, + 1, + 0, + 0, + gather_params, + nullptr), + "cuLaunchKernel(split warmup gather)"); + throw_if_bad(cuCtxSynchronize(), "cuCtxSynchronize(split warmup)"); + } + + std::vector total_timings; + std::vector compute_timings; + std::vector gather_timings; + total_timings.reserve(static_cast(iterations)); + compute_timings.reserve(static_cast(iterations)); + gather_timings.reserve(static_cast(iterations)); + for (int iteration = 0; iteration < iterations; ++iteration) { + throw_if_bad(cuEventRecord(start_event.value, 0), "cuEventRecord(split start)"); + throw_if_bad(cuLaunchKernel(compute_function, + static_cast(blocks), + 1, + 1, + static_cast(threads), + 1, + 1, + 0, + 0, + compute_params, + nullptr), + "cuLaunchKernel(split compute)"); + throw_if_bad(cuEventRecord(compute_stop_event.value, 0), "cuEventRecord(split compute stop)"); + throw_if_bad(cuLaunchKernel(gather_function, + static_cast(blocks), + 1, + 1, + static_cast(threads), + 1, + 1, + 0, + 0, + gather_params, + nullptr), + "cuLaunchKernel(split gather)"); + throw_if_bad(cuEventRecord(stop_event.value, 0), "cuEventRecord(split stop)"); + throw_if_bad(cuEventSynchronize(stop_event.value), "cuEventSynchronize(split stop)"); + + float total_ms = 0.0f; + float compute_ms = 0.0f; + float gather_ms = 0.0f; + throw_if_bad(cuEventElapsedTime(&total_ms, start_event.value, stop_event.value), + "cuEventElapsedTime(split total)"); + throw_if_bad(cuEventElapsedTime(&compute_ms, start_event.value, compute_stop_event.value), + "cuEventElapsedTime(split compute)"); + throw_if_bad(cuEventElapsedTime(&gather_ms, compute_stop_event.value, stop_event.value), + "cuEventElapsedTime(split gather)"); + total_timings.push_back(static_cast(total_ms)); + compute_timings.push_back(static_cast(compute_ms)); + gather_timings.push_back(static_cast(gather_ms)); + } + + split_timing_stats stats; + stats.total = summarize_timings(total_timings); + stats.compute = summarize_timings(compute_timings); + stats.gather = summarize_timings(gather_timings); + return stats; +} + +timing_stats summarize_timings(std::vector const& timings) +{ + timing_stats stats; + if (timings.empty()) { return stats; } + + auto const [min_it, max_it] = std::minmax_element(timings.begin(), timings.end()); + double total = 0.0; + for (double value : timings) { total += value; } + stats.avg = total / static_cast(timings.size()); + stats.min = *min_it; + stats.max = *max_it; + return stats; +} + +void run_cpu_bulk_index_once(PH_ONLINE_JIT_TABLE* table, + std::vector const& keys32, + int vector_width, + std::vector* indexes) +{ + if (indexes == nullptr) { + throw std::runtime_error("indexes output cannot be null"); + } + indexes->resize(keys32.size()); + + std::size_t index = 0; + auto* output = indexes->data(); + auto const* input = keys32.data(); + + auto call_scalar = [&](std::size_t i) { + throw_if_bad(PhOnlineJitIndex32(table, input[i], &output[i]), "PhOnlineJitIndex32"); + }; + + if (vector_width == 16) { + for (; index + 16 <= keys32.size(); index += 16) { + throw_if_bad(PhOnlineJitIndex32x16(table, input + index, output + index), + "PhOnlineJitIndex32x16"); + } + } else if (vector_width == 8) { + for (; index + 8 <= keys32.size(); index += 8) { + throw_if_bad(PhOnlineJitIndex32x8(table, input + index, output + index), + "PhOnlineJitIndex32x8"); + } + } else if (vector_width == 4) { + for (; index + 4 <= keys32.size(); index += 4) { + throw_if_bad(PhOnlineJitIndex32x4(table, input + index, output + index), + "PhOnlineJitIndex32x4"); + } + } else if (vector_width == 2) { + for (; index + 2 <= keys32.size(); index += 2) { + throw_if_bad(PhOnlineJitIndex32x2(table, input + index, output + index), + "PhOnlineJitIndex32x2"); + } + } + + for (; index < keys32.size(); ++index) { + call_scalar(index); + } +} + +timing_stats benchmark_cpu_lookup(PH_ONLINE_JIT_TABLE* table, + std::vector const& keys32, + int vector_width, + int warmup, + int iterations, + std::vector* last_indexes) +{ + for (int iteration = 0; iteration < warmup; ++iteration) { + run_cpu_bulk_index_once(table, keys32, vector_width, last_indexes); + } + + std::vector timings; + timings.reserve(static_cast(iterations)); + for (int iteration = 0; iteration < iterations; ++iteration) { + auto const start = steady_clock::now(); + run_cpu_bulk_index_once(table, keys32, vector_width, last_indexes); + auto const stop = steady_clock::now(); + timings.push_back(elapsed_ms(start, stop)); + } + + return summarize_timings(timings); +} + +void verify_indexes(std::vector const& indexes, + std::vector const& build_keys, + std::vector const& probe_keys, + benchmark_result* result) +{ + if (indexes.empty()) { return; } + if (indexes.size() != probe_keys.size()) { + throw std::runtime_error("Verification indexes must match probe stream length"); + } + + auto const [min_it, max_it] = std::minmax_element(indexes.begin(), indexes.end()); + result->index_min = *min_it; + result->index_max = *max_it; + result->index_span = static_cast(result->index_max) + 1; + + for (std::size_t i = 0; i < indexes.size(); ++i) { + auto const index = indexes[i]; + if (static_cast(index) >= build_keys.size()) { + throw std::runtime_error("Out-of-range index detected at position " + std::to_string(i)); + } + if (build_keys[static_cast(index)] != probe_keys[i]) { + throw std::runtime_error("Key mismatch detected at position " + std::to_string(i)); + } + } +} + +void print_csv_header() +{ + std::cout + << "key_source,key_count,key_bytes,probe_source,probe_key_count,probe_key_bytes,table_data_bytes,table_data_elements," + "table_data_element_size,embed_table_data,hash,mode,lookup_mode,table_load_mode,device_ordinal,device_name," + "arch,items_per_thread,threads,blocks,warmup,iterations,cpu_backend_req," + "cpu_backend_eff,cpu_vector_req,cpu_vector_eff,cpu_strict,cpu_key_mode," + "cpu_compile_hr,cpu_compile_ms,cpu_lookup_avg_ms,cpu_lookup_min_ms," + "cpu_lookup_max_ms,cpu_lookup_ns_per_key,fragment_bytes,combined_bytes,image_bytes," + "host_rss_before_build_bytes,host_rss_after_build_bytes,host_rss_delta_bytes,host_peak_rss_bytes," + "vram_explicit_bytes,vram_total_bytes,build_ms,table_export_ms,emit_ms,compose_ms,compile_ms,link_ms,module_load_ms,alloc_ms,h2d_ms,table_h2d_ms," + "slot_compute_avg_ms,slot_gather_avg_ms,kernel_avg_ms,kernel_min_ms,kernel_max_ms,d2h_ms,verify_ms,index_min," + "index_max,index_span,slot_compute_ns_per_key,slot_gather_ns_per_key,gpu_lookup_ns_per_key," + "warp_unique_requests_avg,warp_duplicate_ratio,block_unique_requests_avg,block_duplicate_ratio\n"; +} + +void print_csv_row(benchmark_result const& result) +{ + std::cout << csv_escape(result.key_source) << ',' + << result.key_count << ',' + << result.key_bytes << ',' + << csv_escape(result.probe_source) << ',' + << result.probe_key_count << ',' + << result.probe_key_bytes << ',' + << result.table_data_bytes << ',' + << result.table_data_elements << ',' + << result.table_data_element_size << ',' + << (result.embedded_table_data ? 1 : 0) << ',' + << csv_escape(result.hash_name) << ',' + << csv_escape(compile_mode_to_string(result.mode)) << ',' + << csv_escape(lookup_mode_to_string(result.lookup)) << ',' + << csv_escape(table_load_mode_to_string(result.table_load)) << ',' + << result.device_ordinal << ',' + << csv_escape(result.device_name) << ',' + << csv_escape("compute_" + std::to_string(result.major) + + std::to_string(result.minor)) << ',' + << result.items_per_thread << ',' + << result.threads << ',' + << result.blocks << ',' + << result.warmup << ',' + << result.iterations << ',' + << csv_escape(result.cpu_backend_requested) << ',' + << csv_escape(result.cpu_backend_effective) << ',' + << result.cpu_vector_requested << ',' + << result.cpu_vector_effective << ',' + << (result.cpu_strict_vector_width ? 1 : 0) << ',' + << csv_escape(result.cpu_key_mode) << ',' + << csv_escape(to_hex(static_cast(result.cpu_compile_hr))) << ',' + << result.cpu_compile_ms << ',' + << result.cpu_lookup_avg_ms << ',' + << result.cpu_lookup_min_ms << ',' + << result.cpu_lookup_max_ms << ',' + << result.cpu_lookup_ns_per_key << ',' + << result.fragment_bytes << ',' + << result.combined_bytes << ',' + << result.image_bytes << ',' + << result.host_rss_before_build_bytes << ',' + << result.host_rss_after_build_bytes << ',' + << result.host_rss_delta_bytes << ',' + << result.host_peak_rss_bytes << ',' + << result.vram_explicit_bytes << ',' + << result.vram_total_bytes << ',' + << std::fixed << std::setprecision(3) + << result.build_ms << ',' + << result.table_export_ms << ',' + << result.emit_ms << ',' + << result.compose_ms << ',' + << result.compile_ms << ',' + << result.link_ms << ',' + << result.module_load_ms << ',' + << result.alloc_ms << ',' + << result.h2d_ms << ',' + << result.table_h2d_ms << ',' + << result.slot_compute_avg_ms << ',' + << result.slot_gather_avg_ms << ',' + << result.kernel_avg_ms << ',' + << result.kernel_min_ms << ',' + << result.kernel_max_ms << ',' + << result.d2h_ms << ',' + << result.verify_ms << ',' + << result.index_min << ',' + << result.index_max << ',' + << result.index_span << ',' + << result.slot_compute_ns_per_key << ',' + << result.slot_gather_ns_per_key << ',' + << result.gpu_lookup_ns_per_key << ',' + << result.warp_unique_requests_avg << ',' + << result.warp_duplicate_ratio << ',' + << result.block_unique_requests_avg << ',' + << result.block_duplicate_ratio << '\n'; +} + +void print_human_summary(benchmark_result const& result) +{ + std::cout << "Keys: " << result.key_count << " (" << format_bytes(result.key_bytes) + << ") from " << result.key_source << "\n"; + std::cout << "Probe: " << result.probe_key_count << " (" << format_bytes(result.probe_key_bytes) + << ") from " << result.probe_source << "\n"; + std::cout << "Table data: " << result.table_data_elements + << " elements x " << result.table_data_element_size + << " bytes = " << format_bytes(result.table_data_bytes) + << ", embedded=" << (result.embedded_table_data ? "yes" : "no") << "\n"; + std::cout << "Config: hash=" << result.hash_name + << ", mode=" << compile_mode_to_string(result.mode) + << ", lookup=" << lookup_mode_to_string(result.lookup) + << ", table-load=" << table_load_mode_to_string(result.table_load) + << ", device=" << result.device_ordinal << " (" << result.device_name << ")" + << ", arch=compute_" << result.major << result.minor + << ", items/thread=" << result.items_per_thread + << ", threads/block=" << result.threads + << ", blocks=" << result.blocks + << ", warmup=" << result.warmup + << ", iterations=" << result.iterations << "\n"; + if (result.cpu_enabled) { + std::cout << "CPU: backend=" << result.cpu_backend_requested + << " -> " << result.cpu_backend_effective + << ", vector=" << result.cpu_vector_requested + << " -> " << result.cpu_vector_effective + << ", strict=" << (result.cpu_strict_vector_width ? 1 : 0) + << ", key-mode=" << result.cpu_key_mode << "\n"; + } + std::cout << "Sizes: fragment=" << result.fragment_bytes + << " bytes, combined=" << result.combined_bytes + << " bytes, image=" << result.image_bytes << " bytes\n"; + std::cout << "Memory: host_rss_before_build=" << result.host_rss_before_build_bytes + << " bytes, host_rss_after_build=" << result.host_rss_after_build_bytes + << " bytes, host_rss_delta=" << result.host_rss_delta_bytes + << " bytes, host_peak_rss=" << result.host_peak_rss_bytes + << " bytes, vram_explicit=" << result.vram_explicit_bytes + << " bytes, vram_total=" << result.vram_total_bytes << " bytes\n"; + std::cout << std::fixed << std::setprecision(3) + << "Timings (ms): build=" << result.build_ms + << ", table_export=" << result.table_export_ms + << ", emit=" << result.emit_ms + << ", compose=" << result.compose_ms + << ", cpu_compile=" << result.cpu_compile_ms + << ", cpu_lookup_avg=" << result.cpu_lookup_avg_ms + << ", compile=" << result.compile_ms + << ", link=" << result.link_ms + << ", module_load=" << result.module_load_ms + << ", alloc=" << result.alloc_ms + << ", h2d=" << result.h2d_ms + << ", table_h2d=" << result.table_h2d_ms + << ", slot_compute=" << result.slot_compute_avg_ms + << ", slot_gather=" << result.slot_gather_avg_ms + << ", kernel_avg=" << result.kernel_avg_ms + << ", kernel_min=" << result.kernel_min_ms + << ", kernel_max=" << result.kernel_max_ms + << ", d2h=" << result.d2h_ms + << ", verify=" << result.verify_ms << "\n"; + std::cout << "Throughput: gpu=" << std::setprecision(3) << result.gpu_lookup_ns_per_key + << " ns/key"; + if (result.lookup == lookup_mode::split) { + std::cout << " (slot_compute=" << result.slot_compute_ns_per_key + << ", slot_gather=" << result.slot_gather_ns_per_key << ")"; + } + if (result.cpu_enabled) { + std::cout << ", cpu=" << result.cpu_lookup_ns_per_key << " ns/key"; + } + std::cout << "\n"; + if (result.slot_reuse_analyzed) { + std::cout << "Slot reuse: warp unique avg=" << result.warp_unique_requests_avg + << ", warp duplicate ratio=" << result.warp_duplicate_ratio + << ", block unique avg=" << result.block_unique_requests_avg + << ", block duplicate ratio=" << result.block_duplicate_ratio << "\n"; + } + if (result.index_span != 0) { + double const load_factor = + static_cast(result.key_count) / static_cast(result.index_span); + std::cout << "Index range: min=" << result.index_min + << ", max=" << result.index_max + << ", span=" << result.index_span + << ", load-factor=" << std::setprecision(6) << load_factor << "\n"; + } +} + +options parse_args(int argc, char** argv) +{ + options opts; + + for (int index = 1; index < argc; ++index) { + std::string arg = argv[index]; + if (arg == "--help" || arg == "-h") { + print_usage(argv[0]); + std::exit(0); + } + if (arg == "--hash" && index + 1 < argc) { + opts.hash_name = argv[++index]; + continue; + } + if (arg == "--items-per-thread" && index + 1 < argc) { + opts.items_per_thread = std::stoi(argv[++index]); + continue; + } + if (arg == "--threads" && index + 1 < argc) { + opts.threads = std::stoi(argv[++index]); + continue; + } + if (arg == "--iterations" && index + 1 < argc) { + opts.iterations = std::stoi(argv[++index]); + continue; + } + if (arg == "--warmup" && index + 1 < argc) { + opts.warmup = std::stoi(argv[++index]); + continue; + } + if (arg == "--compile-mode" && index + 1 < argc) { + opts.compile_mode_name = argv[++index]; + continue; + } + if (arg == "--lookup-mode" && index + 1 < argc) { + opts.lookup_mode_name = argv[++index]; + continue; + } + if (arg == "--table-load-mode" && index + 1 < argc) { + opts.table_load_mode_name = argv[++index]; + continue; + } + if (arg == "--cpu-backend" && index + 1 < argc) { + opts.cpu_backend_name = argv[++index]; + continue; + } + if (arg == "--cpu-vector-width" && index + 1 < argc) { + opts.cpu_vector_width = std::stoi(argv[++index]); + continue; + } + if (arg == "--cpu-strict-vector-width" && index + 1 < argc) { + opts.cpu_strict_vector_width = (std::stoi(argv[++index]) != 0); + continue; + } + if (arg == "--analyze-slot-reuse") { + opts.analyze_slot_reuse = true; + continue; + } + if (arg == "--analysis-only") { + opts.analyze_slot_reuse = true; + opts.analysis_only = true; + continue; + } + if (arg == "--device" && index + 1 < argc) { + opts.device_ordinal = std::stoi(argv[++index]); + continue; + } + if (arg == "--keys-file" && index + 1 < argc) { + opts.keys_file = argv[++index]; + continue; + } + if (arg == "--probe-keys-file" && index + 1 < argc) { + opts.probe_keys_file = argv[++index]; + continue; + } + if (arg == "--max-keys" && index + 1 < argc) { + opts.max_keys = std::stoull(argv[++index]); + continue; + } + if (arg == "--max-probe-keys" && index + 1 < argc) { + opts.max_probe_keys = std::stoull(argv[++index]); + continue; + } + if (arg == "--dump-fragment") { + opts.dump_fragment = true; + continue; + } + if (arg == "--source-out" && index + 1 < argc) { + opts.source_out_path = argv[++index]; + continue; + } + if (arg == "--csv") { + opts.csv = true; + continue; + } + if (arg == "--csv-header") { + opts.csv = true; + opts.csv_header = true; + continue; + } + if (arg == "--embed-table-data") { + opts.embed_table_data = true; + continue; + } + if (arg == "--no-verify") { + opts.verify = false; + continue; + } + + std::ostringstream message; + message << "Unknown argument: " << arg; + throw std::runtime_error(message.str()); + } + + if (!is_supported_hash_name(opts.hash_name)) { + throw std::runtime_error("Unsupported hash: " + opts.hash_name); + } + if (opts.compile_mode_name != "ptx" && opts.compile_mode_name != "lto") { + throw std::runtime_error("Unsupported compile mode: " + opts.compile_mode_name); + } + if (opts.lookup_mode_name != "direct" && + opts.lookup_mode_name != "split" && + opts.lookup_mode_name != "warpcache" && + opts.lookup_mode_name != "blocksort") { + throw std::runtime_error("Unsupported lookup mode: " + opts.lookup_mode_name); + } + if (opts.table_load_mode_name != "generic" && + opts.table_load_mode_name != "readonly") { + throw std::runtime_error("Unsupported table-load mode: " + opts.table_load_mode_name); + } + if (!is_supported_backend_name(opts.cpu_backend_name)) { + throw std::runtime_error("Unsupported CPU backend: " + opts.cpu_backend_name); + } + if (opts.cpu_vector_width != 1 && opts.cpu_vector_width != 2 && + opts.cpu_vector_width != 4 && opts.cpu_vector_width != 8 && + opts.cpu_vector_width != 16) { + throw std::runtime_error("cpu-vector-width must be one of 1,2,4,8,16"); + } + if (opts.items_per_thread <= 0) { + throw std::runtime_error("items-per-thread must be > 0"); + } + if (opts.threads <= 0) { + throw std::runtime_error("threads must be > 0"); + } + if (opts.iterations <= 0) { + throw std::runtime_error("iterations must be > 0"); + } + if (opts.warmup < 0) { + throw std::runtime_error("warmup must be >= 0"); + } + + return opts; +} + +benchmark_result run_benchmark(options const& opts) +{ + benchmark_result result; + auto const selected_hash = parse_hash_function(opts.hash_name); + result.hash_name = opts.hash_name; + result.mode = parse_compile_mode(opts.compile_mode_name); + result.lookup = parse_lookup_mode(opts.lookup_mode_name); + result.table_load = parse_table_load_mode(opts.table_load_mode_name); + result.device_ordinal = opts.device_ordinal; + result.items_per_thread = opts.items_per_thread; + result.threads = opts.threads; + result.warmup = opts.warmup; + result.iterations = opts.iterations; + result.embedded_table_data = opts.embed_table_data; + result.cpu_enabled = (opts.cpu_backend_name != "none"); + result.cpu_backend_requested = opts.cpu_backend_name; + result.cpu_vector_requested = opts.cpu_vector_width; + result.cpu_strict_vector_width = opts.cpu_strict_vector_width; + if (result.cpu_enabled && opts.cpu_backend_name == "rawdog-jit" && opts.cpu_vector_width == 2) { + throw std::runtime_error("cpu-backend=rawdog-jit does not support cpu-vector-width=2"); + } + bool const validate_membership = !opts.probe_keys_file.empty(); + + if (validate_membership && result.lookup != lookup_mode::direct) { + throw std::runtime_error( + "Separate build/probe streams are currently only supported for lookup-mode=direct"); + } + if (opts.verify && result.lookup != lookup_mode::direct) { + throw std::runtime_error("--verify is currently only supported for lookup-mode=direct"); + } + + if (result.lookup == lookup_mode::blocksort) { + auto const total_requests = + static_cast(opts.threads) * + static_cast(opts.items_per_thread) * 2u; + if (total_requests == 0 || (total_requests & (total_requests - 1u)) != 0) { + throw std::runtime_error( + "blocksort lookup mode requires threads * items_per_thread * 2 to be a power of two"); + } + } + + auto build_keys = opts.keys_file.empty() ? make_sample_keys() + : load_keys_file(opts.keys_file, opts.max_keys); + if (build_keys.empty()) { + throw std::runtime_error("No keys available to benchmark"); + } + { + auto sorted_build_keys = build_keys; + std::sort(sorted_build_keys.begin(), sorted_build_keys.end()); + if (std::adjacent_find(sorted_build_keys.begin(), sorted_build_keys.end()) != + sorted_build_keys.end()) { + throw std::runtime_error("Build keys must be unique for the PerfectHash benchmark"); + } + } + result.key_source = opts.keys_file.empty() ? std::string{"built-in-sample"} : opts.keys_file; + result.key_count = build_keys.size(); + result.key_bytes = build_keys.size() * sizeof(std::uint64_t); + auto probe_keys = opts.probe_keys_file.empty() + ? build_keys + : load_keys_file(opts.probe_keys_file, opts.max_probe_keys); + if (opts.probe_keys_file.empty() && opts.max_probe_keys > 0 && + probe_keys.size() > opts.max_probe_keys) { + probe_keys.resize(static_cast(opts.max_probe_keys)); + } + if (probe_keys.empty()) { + throw std::runtime_error("No probe keys available to benchmark"); + } + result.probe_source = opts.probe_keys_file.empty() ? result.key_source : opts.probe_keys_file; + result.probe_key_count = probe_keys.size(); + result.probe_key_bytes = probe_keys.size() * sizeof(std::uint64_t); + result.host_rss_before_build_bytes = current_rss_bytes(); + + ph_context_handle context; + ph_table_handle table; + cuda_source_handle fragment; + table_data_handle table_data; + PH_ONLINE_JIT_TABLE_INFO cpu_table_info{}; + + auto build_start = steady_clock::now(); + throw_if_bad(PhOnlineJitOpen(&context.value), "PhOnlineJitOpen"); + throw_if_bad(PhOnlineJitCreateTable64(context.value, + selected_hash, + build_keys.data(), + static_cast(build_keys.size()), + &table.value), + "PhOnlineJitCreateTable64"); + throw_if_bad(PhOnlineJitGetTableInfo(table.value, &cpu_table_info), + "PhOnlineJitGetTableInfo"); + auto const table_hash = + static_cast(cpu_table_info.HashFunctionId); + auto build_stop = steady_clock::now(); + result.build_ms = elapsed_ms(build_start, build_stop); + result.host_rss_after_build_bytes = current_rss_bytes(); + result.host_rss_delta_bytes = + (result.host_rss_after_build_bytes >= result.host_rss_before_build_bytes) + ? (result.host_rss_after_build_bytes - result.host_rss_before_build_bytes) + : 0; + result.host_peak_rss_bytes = peak_rss_bytes(); + + if (opts.analyze_slot_reuse) { + auto const warp_stats = + analyze_slot_reuse_for_tile(probe_keys, table_hash, cpu_table_info, 32u, static_cast(opts.items_per_thread)); + auto const block_stats = + analyze_slot_reuse_for_tile(probe_keys, + table_hash, + cpu_table_info, + static_cast(opts.threads), + static_cast(opts.items_per_thread)); + result.slot_reuse_analyzed = true; + result.warp_unique_requests_avg = warp_stats.unique_requests_avg; + result.warp_duplicate_ratio = warp_stats.duplicate_ratio; + result.block_unique_requests_avg = block_stats.unique_requests_avg; + result.block_duplicate_ratio = block_stats.duplicate_ratio; + } + + if (opts.analysis_only) { + result.cpu_enabled = false; + return result; + } + + if (result.cpu_enabled) { + if (cpu_table_info.KeySizeInBytes > 4) { + throw std::runtime_error( + "CPU bulk-index benchmark requires a table downsized to 32-bit keys"); + } + if (cpu_table_info.OriginalKeySizeInBytes > 4) { + if (cpu_table_info.DownsizeBitmap == 0) { + throw std::runtime_error( + "CPU bulk-index benchmark requires a lossless 64-to-32 downsizing transform"); + } + if (probe_keys != build_keys) { + throw std::runtime_error( + "CPU bulk-index benchmark on downsized 64-bit tables is only supported for self-probe workloads"); + } + } + + preload_llvm_runtime_library(opts.cpu_backend_name); + + PH_ONLINE_JIT_BACKEND effective_backend = PhOnlineJitBackendAuto; + std::uint32_t effective_vector_width = 0; + std::uint32_t compile_flags = + opts.cpu_strict_vector_width ? PH_ONLINE_JIT_COMPILE_FLAG_STRICT_VECTOR_WIDTH : 0; + + auto cpu_compile_start = steady_clock::now(); + result.cpu_compile_hr = PhOnlineJitCompileTableEx(context.value, + table.value, + parse_backend(opts.cpu_backend_name), + static_cast(opts.cpu_vector_width), + PhOnlineJitMaxIsaAuto, + compile_flags, + &effective_backend, + &effective_vector_width); + auto cpu_compile_stop = steady_clock::now(); + result.cpu_compile_ms = elapsed_ms(cpu_compile_start, cpu_compile_stop); + throw_if_bad(result.cpu_compile_hr, "PhOnlineJitCompileTableEx"); + + result.cpu_jit_enabled = true; + result.cpu_backend_effective = backend_to_string(effective_backend); + result.cpu_vector_effective = static_cast(effective_vector_width == 0 ? 1 : effective_vector_width); + result.cpu_key_mode = "downsized32"; + + auto const downsized_keys32 = make_downsized_keys32(probe_keys, cpu_table_info); + std::vector cpu_indexes; + auto const cpu_stats = + benchmark_cpu_lookup(table.value, + downsized_keys32, + result.cpu_vector_effective, + opts.warmup, + opts.iterations, + &cpu_indexes); + result.cpu_lookup_avg_ms = cpu_stats.avg; + result.cpu_lookup_min_ms = cpu_stats.min; + result.cpu_lookup_max_ms = cpu_stats.max; + result.cpu_lookup_ns_per_key = ns_per_key(result.cpu_lookup_avg_ms, result.probe_key_count); + } + + if (!opts.embed_table_data || opts.verify) { + auto table_export_start = steady_clock::now(); + throw_if_bad(PhOnlineJitGetCudaTableData(table.value, + &table_data.value, + &result.table_data_bytes, + &result.table_data_element_size, + &result.table_data_elements), + "PhOnlineJitGetCudaTableData"); + auto table_export_stop = steady_clock::now(); + result.table_export_ms = elapsed_ms(table_export_start, table_export_stop); + } else { + result.table_data_bytes = + static_cast(cpu_table_info.NumberOfTableElements) * + static_cast(cpu_table_info.AssignedElementSizeInBytes); + result.table_data_elements = static_cast(cpu_table_info.NumberOfTableElements); + result.table_data_element_size = cpu_table_info.AssignedElementSizeInBytes; + } + + auto emit_start = steady_clock::now(); + auto source_flags = PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_KERNELS; + if (!opts.embed_table_data) { + source_flags |= PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA; + } + throw_if_bad(PhOnlineJitGetCudaSourceEx(table.value, + source_flags, + &fragment.value, + &result.fragment_bytes), + "PhOnlineJitGetCudaSourceEx"); + auto emit_stop = steady_clock::now(); + result.emit_ms = elapsed_ms(emit_start, emit_stop); + + std::string fragment_source(fragment.value, result.fragment_bytes); + if (opts.dump_fragment) { std::cout << fragment_source << "\n"; } + + auto compose_start = steady_clock::now(); + std::string combined_source = + compose_translation_unit(fragment_source, + opts.items_per_thread, + opts.threads, + opts.embed_table_data, + validate_membership, + result.lookup, + result.table_load); + auto compose_stop = steady_clock::now(); + result.compose_ms = elapsed_ms(compose_start, compose_stop); + result.combined_bytes = combined_source.size(); + + if (!opts.source_out_path.empty()) { + std::ofstream out(opts.source_out_path, std::ios::binary); + out.write(combined_source.data(), + static_cast(combined_source.size())); + } + + throw_if_bad(cuInit(0), "cuInit"); + + CUdevice device{}; + throw_if_bad(cuDeviceGet(&device, opts.device_ordinal), "cuDeviceGet"); + char device_name[128] = {}; + throw_if_bad(cuDeviceGetName(device_name, static_cast(sizeof(device_name)), device), + "cuDeviceGetName"); + result.device_name = device_name; + throw_if_bad(cuDeviceGetAttribute(&result.major, + CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + device), + "cuDeviceGetAttribute(major)"); + throw_if_bad(cuDeviceGetAttribute(&result.minor, + CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, + device), + "cuDeviceGetAttribute(minor)"); + if (result.lookup == lookup_mode::warpcache && result.major < 7) { + throw std::runtime_error("lookup-mode=warpcache requires compute capability 7.0 or newer"); + } + if (result.lookup == lookup_mode::blocksort) { + int shared_limit = 0; + auto const shared_bytes = + static_cast(opts.threads) * + static_cast(opts.items_per_thread) * 2u * + sizeof(std::uint32_t) * 4u; + throw_if_bad(cuDeviceGetAttribute(&shared_limit, + CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK, + device), + "cuDeviceGetAttribute(max_shared_memory_per_block)"); + if (shared_bytes > static_cast(shared_limit)) { + throw std::runtime_error( + "blocksort lookup mode exceeds the device shared-memory-per-block limit"); + } + } + + primary_context_handle primary_context; + primary_context.device = device; + CUcontext cu_context{}; + throw_if_bad(cuDevicePrimaryCtxRetain(&cu_context, device), "cuDevicePrimaryCtxRetain"); + primary_context.retained = true; + throw_if_bad(cuCtxSetCurrent(cu_context), "cuCtxSetCurrent"); + + std::size_t free_before_gpu = 0; + std::size_t total_gpu_mem = 0; + throw_if_bad(cuCtxSynchronize(), "cuCtxSynchronize(before meminfo)"); + throw_if_bad(cudaMemGetInfo(&free_before_gpu, &total_gpu_mem), "cudaMemGetInfo(before)"); + + std::string ptx; + std::vector cubin; + if (result.mode == compile_mode::ptx) { + auto compile_start = steady_clock::now(); + ptx = compile_to_ptx(combined_source, result.major, result.minor); + auto compile_stop = steady_clock::now(); + result.compile_ms = elapsed_ms(compile_start, compile_stop); + result.image_bytes = ptx.size(); + } else { + auto compile_start = steady_clock::now(); + auto ir = compile_to_ltoir(combined_source, result.major, result.minor); + auto compile_stop = steady_clock::now(); + result.compile_ms = elapsed_ms(compile_start, compile_stop); + + auto link_start = steady_clock::now(); + cubin = link_ltoir_to_cubin(ir, result.major, result.minor); + auto link_stop = steady_clock::now(); + result.link_ms = elapsed_ms(link_start, link_stop); + result.image_bytes = cubin.size(); + } + + module_handle module; + void const* image = (result.mode == compile_mode::ptx) + ? static_cast(ptx.data()) + : static_cast(cubin.data()); + auto module_load_start = steady_clock::now(); + throw_if_bad(cuModuleLoadDataEx(&module.value, image, 0, nullptr, nullptr), + "cuModuleLoadDataEx"); + auto module_load_stop = steady_clock::now(); + result.module_load_ms = elapsed_ms(module_load_start, module_load_stop); + + CUfunction direct_function{}; + CUfunction compute_slots_function{}; + CUfunction gather_function{}; + CUfunction warpcache_function{}; + CUfunction blocksort_function{}; + if (result.lookup == lookup_mode::direct) { + throw_if_bad(cuModuleGetFunction(&direct_function, module.value, "probe_kernel"), + "cuModuleGetFunction(probe_kernel)"); + } else if (result.lookup == lookup_mode::split) { + throw_if_bad(cuModuleGetFunction(&compute_slots_function, module.value, "compute_slots_kernel"), + "cuModuleGetFunction(compute_slots_kernel)"); + throw_if_bad(cuModuleGetFunction(&gather_function, module.value, "gather_kernel"), + "cuModuleGetFunction(gather_kernel)"); + } else if (result.lookup == lookup_mode::warpcache) { + throw_if_bad(cuModuleGetFunction(&warpcache_function, module.value, "warpcache_probe_kernel"), + "cuModuleGetFunction(warpcache_probe_kernel)"); + } else { + throw_if_bad(cuModuleGetFunction(&blocksort_function, module.value, "blocksort_probe_kernel"), + "cuModuleGetFunction(blocksort_probe_kernel)"); + } + + device_memory d_keys; + device_memory d_build_keys; + device_memory d_table; + device_memory d_slot1; + device_memory d_slot2; + device_memory d_indexes; + auto alloc_start = steady_clock::now(); + throw_if_bad(cuMemAlloc(&d_keys.value, result.probe_key_bytes), "cuMemAlloc(keys)"); + if (validate_membership) { + throw_if_bad(cuMemAlloc(&d_build_keys.value, result.key_bytes), "cuMemAlloc(build_keys)"); + } + if (!opts.embed_table_data) { + throw_if_bad(cuMemAlloc(&d_table.value, result.table_data_bytes), "cuMemAlloc(table)"); + } + if (result.lookup == lookup_mode::split) { + throw_if_bad(cuMemAlloc(&d_slot1.value, + probe_keys.size() * sizeof(std::uint32_t)), + "cuMemAlloc(slot1)"); + throw_if_bad(cuMemAlloc(&d_slot2.value, + probe_keys.size() * sizeof(std::uint32_t)), + "cuMemAlloc(slot2)"); + } + throw_if_bad(cuMemAlloc(&d_indexes.value, + probe_keys.size() * sizeof(std::uint32_t)), + "cuMemAlloc(indexes)"); + auto alloc_stop = steady_clock::now(); + result.alloc_ms = elapsed_ms(alloc_start, alloc_stop); + result.vram_explicit_bytes = result.probe_key_bytes + + (validate_membership ? result.key_bytes : 0) + + (!opts.embed_table_data ? result.table_data_bytes : 0) + + (probe_keys.size() * sizeof(std::uint32_t)) + + ((result.lookup == lookup_mode::split) + ? (probe_keys.size() * sizeof(std::uint32_t) * 2) + : 0); + + auto h2d_start = steady_clock::now(); + throw_if_bad(cuMemcpyHtoD(d_keys.value, probe_keys.data(), result.probe_key_bytes), "cuMemcpyHtoD"); + if (validate_membership) { + throw_if_bad(cuMemcpyHtoD(d_build_keys.value, build_keys.data(), result.key_bytes), + "cuMemcpyHtoD(build_keys)"); + } + if (!opts.embed_table_data) { + auto table_h2d_start = steady_clock::now(); + throw_if_bad(cuMemcpyHtoD(d_table.value, + table_data.value, + result.table_data_bytes), + "cuMemcpyHtoD(table)"); + auto table_h2d_stop = steady_clock::now(); + result.table_h2d_ms = elapsed_ms(table_h2d_start, table_h2d_stop); + } + auto h2d_stop = steady_clock::now(); + result.h2d_ms = elapsed_ms(h2d_start, h2d_stop); + + result.blocks = static_cast((probe_keys.size() + + static_cast(opts.threads * opts.items_per_thread) - 1) / + static_cast(opts.threads * opts.items_per_thread)); + std::size_t count = probe_keys.size(); + std::size_t build_key_count = build_keys.size(); + if (result.lookup == lookup_mode::direct) { + void* embedded_params_no_check[] = {&d_keys.value, &d_indexes.value, &count}; + void* embedded_params_check[] = { + &d_keys.value, &d_build_keys.value, &build_key_count, &d_indexes.value, &count}; + void* detached_params_no_check[] = {&d_keys.value, &d_table.value, &d_indexes.value, &count}; + void* detached_params_check[] = { + &d_keys.value, &d_build_keys.value, &build_key_count, &d_table.value, &d_indexes.value, &count}; + auto kernel_timings = run_probe_iterations(direct_function, + opts.embed_table_data + ? (validate_membership ? embedded_params_check : embedded_params_no_check) + : (validate_membership ? detached_params_check : detached_params_no_check), + result.blocks, + opts.threads, + opts.warmup, + opts.iterations); + auto const stats = summarize_timings(kernel_timings); + result.kernel_avg_ms = stats.avg; + result.kernel_min_ms = stats.min; + result.kernel_max_ms = stats.max; + } else if (result.lookup == lookup_mode::split) { + void* compute_params[] = {&d_keys.value, &d_slot1.value, &d_slot2.value, &count}; + if (opts.embed_table_data) { + void* gather_params[] = {&d_slot1.value, &d_slot2.value, &d_indexes.value, &count}; + auto const split_stats = run_split_probe_iterations(compute_slots_function, + compute_params, + gather_function, + gather_params, + result.blocks, + opts.threads, + opts.warmup, + opts.iterations); + result.slot_compute_avg_ms = split_stats.compute.avg; + result.slot_gather_avg_ms = split_stats.gather.avg; + result.kernel_avg_ms = split_stats.total.avg; + result.kernel_min_ms = split_stats.total.min; + result.kernel_max_ms = split_stats.total.max; + } else { + void* gather_params[] = {&d_slot1.value, &d_slot2.value, &d_table.value, &d_indexes.value, &count}; + auto const split_stats = run_split_probe_iterations(compute_slots_function, + compute_params, + gather_function, + gather_params, + result.blocks, + opts.threads, + opts.warmup, + opts.iterations); + result.slot_compute_avg_ms = split_stats.compute.avg; + result.slot_gather_avg_ms = split_stats.gather.avg; + result.kernel_avg_ms = split_stats.total.avg; + result.kernel_min_ms = split_stats.total.min; + result.kernel_max_ms = split_stats.total.max; + } + } else if (result.lookup == lookup_mode::warpcache) { + void* embedded_params[] = {&d_keys.value, &d_indexes.value, &count}; + void* detached_params[] = {&d_keys.value, &d_table.value, &d_indexes.value, &count}; + auto kernel_timings = run_probe_iterations(warpcache_function, + opts.embed_table_data ? embedded_params : detached_params, + result.blocks, + opts.threads, + opts.warmup, + opts.iterations); + auto const stats = summarize_timings(kernel_timings); + result.kernel_avg_ms = stats.avg; + result.kernel_min_ms = stats.min; + result.kernel_max_ms = stats.max; + } else { + void* embedded_params[] = {&d_keys.value, &d_indexes.value, &count}; + void* detached_params[] = {&d_keys.value, &d_table.value, &d_indexes.value, &count}; + auto kernel_timings = run_probe_iterations(blocksort_function, + opts.embed_table_data ? embedded_params : detached_params, + result.blocks, + opts.threads, + opts.warmup, + opts.iterations); + auto const stats = summarize_timings(kernel_timings); + result.kernel_avg_ms = stats.avg; + result.kernel_min_ms = stats.min; + result.kernel_max_ms = stats.max; + } + + std::vector indexes(probe_keys.size()); + auto d2h_start = steady_clock::now(); + throw_if_bad(cuMemcpyDtoH(indexes.data(), + d_indexes.value, + indexes.size() * sizeof(std::uint32_t)), + "cuMemcpyDtoH"); + auto d2h_stop = steady_clock::now(); + result.d2h_ms = elapsed_ms(d2h_start, d2h_stop); + + result.slot_compute_ns_per_key = ns_per_key(result.slot_compute_avg_ms, result.probe_key_count); + result.slot_gather_ns_per_key = ns_per_key(result.slot_gather_avg_ms, result.probe_key_count); + result.gpu_lookup_ns_per_key = ns_per_key(result.kernel_avg_ms, result.probe_key_count); + std::size_t free_after_gpu = 0; + throw_if_bad(cuCtxSynchronize(), "cuCtxSynchronize(after meminfo)"); + throw_if_bad(cudaMemGetInfo(&free_after_gpu, &total_gpu_mem), "cudaMemGetInfo(after)"); + result.vram_total_bytes = + (free_before_gpu >= free_after_gpu) ? (free_before_gpu - free_after_gpu) : 0; + + if (opts.verify && result.lookup == lookup_mode::direct) { + auto verify_start = steady_clock::now(); + if (probe_keys == build_keys) { + verify_indexes(indexes, build_keys, probe_keys, &result); + } else { + if (table_data.value == nullptr) { + throw std::runtime_error("Verification requires host-accessible table data"); + } + for (std::size_t i = 0; i < probe_keys.size(); ++i) { + auto const candidate = + index_from_key_host(probe_keys[i], table_hash, cpu_table_info, table_data.value); + auto const expected = + (candidate < build_keys.size() && build_keys[candidate] == probe_keys[i]) + ? candidate + : std::numeric_limits::max(); + if (indexes[i] != expected) { + throw std::runtime_error("Verification failed at probe index " + std::to_string(i)); + } + } + if (!indexes.empty()) { + std::uint32_t min_index = std::numeric_limits::max(); + std::uint32_t max_index = 0; + bool saw_hit = false; + for (auto const index : indexes) { + if (index == std::numeric_limits::max()) { + continue; + } + min_index = std::min(min_index, index); + max_index = std::max(max_index, index); + saw_hit = true; + } + if (saw_hit) { + result.index_min = min_index; + result.index_max = max_index; + result.index_span = static_cast(result.index_max) + 1; + } else { + result.index_min = 0; + result.index_max = 0; + result.index_span = 0; + } + } + } + auto verify_stop = steady_clock::now(); + result.verify_ms = elapsed_ms(verify_start, verify_stop); + } + + if (opts.verify && result.cpu_enabled) { + auto const downsized_keys32 = make_downsized_keys32(probe_keys, cpu_table_info); + std::vector cpu_indexes; + run_cpu_bulk_index_once(table.value, + downsized_keys32, + result.cpu_vector_effective, + &cpu_indexes); + if (validate_membership) { + for (std::size_t i = 0; i < cpu_indexes.size(); ++i) { + auto const candidate = cpu_indexes[i]; + cpu_indexes[i] = + (candidate < build_keys.size() && build_keys[candidate] == probe_keys[i]) + ? candidate + : std::numeric_limits::max(); + } + } + if (cpu_indexes != indexes) { + throw std::runtime_error("CPU and GPU index outputs differed"); + } + } + + return result; +} + +} // namespace + +int main(int argc, char** argv) +{ + try { + auto const opts = parse_args(argc, argv); + auto const result = run_benchmark(opts); + if (opts.csv_header) { print_csv_header(); } + if (opts.csv) { print_csv_row(result); } + else { print_human_summary(result); } + return 0; + } catch (std::exception const& e) { + std::cerr << e.what() << "\n"; + return 1; + } +} diff --git a/examples/cpp-console-online-jit/CMakeLists.txt b/examples/cpp-console-online-jit/CMakeLists.txt index 32808470..f69b863d 100644 --- a/examples/cpp-console-online-jit/CMakeLists.txt +++ b/examples/cpp-console-online-jit/CMakeLists.txt @@ -46,6 +46,10 @@ set(PH_ONLINE_JIT_TARGET "") set(PH_ONLINE_JIT_RUNTIME_TARGET "") set(PH_ONLINE_JIT_LLVM_TARGET "") +if(PERFECTHASH_ROOT AND NOT PERFECTHASH_ROOT STREQUAL "") + set(PH_ONLINE_JIT_USE_FETCHCONTENT OFF CACHE BOOL "" FORCE) +endif() + if(PH_ONLINE_JIT_USE_FETCHCONTENT) include(FetchContent) diff --git a/examples/cpp-console-online-jit/README.md b/examples/cpp-console-online-jit/README.md index 760c56a5..070137fc 100644 --- a/examples/cpp-console-online-jit/README.md +++ b/examples/cpp-console-online-jit/README.md @@ -77,6 +77,24 @@ Optional arguments: - `--backend ` - `--hash ` - `--vector-width <0|1|2|4|8|16>` +- `--dump-cuda-source` +- `--omit-kernels` +- `--source-out ` + +### Dump Generated CUDA Source + +```bash +./build/examples/cpp-console-online-jit/cpp-console-online-jit \ + --backend rawdog-jit \ + --hash mulshrolate3rx \ + --source-out /tmp/online_jit_table.cu \ + --omit-kernels +``` + +`--omit-kernels` emits only the inline lookup fragment (`index_from_key()`, +constants, table data, and helpers) and leaves kernel shape to the downstream +consumer. This is the mode intended for NVRTC/LTO-IR pipelines that want +call-site specialization such as `ITEMS_PER_THREAD` unrolling. ## Notes diff --git a/examples/cpp-console-online-jit/cmake/FindPerfectHashOnlineJit.cmake b/examples/cpp-console-online-jit/cmake/FindPerfectHashOnlineJit.cmake index e2910e5e..fb59b928 100644 --- a/examples/cpp-console-online-jit/cmake/FindPerfectHashOnlineJit.cmake +++ b/examples/cpp-console-online-jit/cmake/FindPerfectHashOnlineJit.cmake @@ -29,10 +29,13 @@ list(REMOVE_DUPLICATES _ph_search_hints) find_path( PERFECTHASH_ONLINE_JIT_INCLUDE_DIR - NAMES PerfectHashOnlineJit.h + NAMES + PerfectHash/PerfectHashOnlineJit.h + PerfectHashOnlineJit.h HINTS ${_ph_search_hints} PATH_SUFFIXES include + . ) set(_ph_lib_suffixes) @@ -105,6 +108,7 @@ endif() find_library( PERFECTHASH_ONLINE_JIT_LIBRARY NAMES + PerfectHashOnlineCore PerfectHashOnline HINTS ${_ph_search_hints} PATH_SUFFIXES ${_ph_lib_suffixes} @@ -140,6 +144,7 @@ if(WIN32) find_file( PERFECTHASH_ONLINE_JIT_RUNTIME_DLL NAMES + PerfectHashOnlineCore.dll PerfectHashOnline.dll HINTS ${_ph_search_hints} diff --git a/examples/cpp-console-online-jit/src/main.cpp b/examples/cpp-console-online-jit/src/main.cpp index 75eb4098..9eabdd2d 100644 --- a/examples/cpp-console-online-jit/src/main.cpp +++ b/examples/cpp-console-online-jit/src/main.cpp @@ -1,5 +1,6 @@ #include +#include #include #include #include @@ -96,7 +97,9 @@ void PreloadLlvmRuntimeLibrary(PH_ONLINE_JIT_BACKEND backend) { void PrintUsage(const char *argv0) { std::cout << "Usage: " << argv0 << " [--backend ]" - " [--hash ] [--vector-width <0|1|2|4|8|16>]\n"; + " [--hash ] [--vector-width <0|1|2|4|8|16>]" + " [--dump-cuda-source] [--omit-kernels]" + " [--source-out ]\n"; std::cout << "Hash names: multiplyshiftr, multiplyshiftrx, mulshrolate1rx,\n"; std::cout << " mulshrolate2rx, mulshrolate3rx, mulshrolate4rx\n"; } @@ -110,7 +113,10 @@ int main(int argc, char **argv) { std::string hash_name = "mulshrolate2rx"; std::string backend_name = "rawdog-jit"; + std::string source_out_path; uint32_t vector_width = 16; + bool dump_cuda_source = false; + bool omit_kernels = false; for (int index = 1; index < argc; ++index) { const std::string arg = argv[index]; @@ -134,6 +140,21 @@ int main(int argc, char **argv) { continue; } + if (arg == "--dump-cuda-source") { + dump_cuda_source = true; + continue; + } + + if (arg == "--omit-kernels") { + omit_kernels = true; + continue; + } + + if (arg == "--source-out" && index + 1 < argc) { + source_out_path = argv[++index]; + continue; + } + std::cerr << "Unknown argument: " << arg << "\n"; PrintUsage(argv[0]); return 2; @@ -177,6 +198,46 @@ int main(int argc, char **argv) { return 1; } + if (dump_cuda_source || !source_out_path.empty()) { + char *cuda_source = nullptr; + size_t cuda_source_size = 0; + uint32_t source_flags = omit_kernels ? PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_KERNELS : 0; + + result = PhOnlineJitGetCudaSourceEx(table, + source_flags, + &cuda_source, + &cuda_source_size); + if (!Succeeded(result)) { + std::cerr << "PhOnlineJitGetCudaSourceEx() failed: " << ToHex(result) + << "\n"; + PhOnlineJitReleaseTable(table); + PhOnlineJitClose(context); + return 1; + } + + if (!source_out_path.empty()) { + std::ofstream out(source_out_path, std::ios::binary); + if (!out) { + std::cerr << "Unable to open source output path: " << source_out_path + << "\n"; + PhOnlineJitFreeCudaSource(cuda_source); + PhOnlineJitReleaseTable(table); + PhOnlineJitClose(context); + return 1; + } + out.write(cuda_source, static_cast(cuda_source_size)); + } + + if (dump_cuda_source) { + std::cout.write(cuda_source, static_cast(cuda_source_size)); + if (cuda_source_size == 0 || cuda_source[cuda_source_size - 1] != '\n') { + std::cout << '\n'; + } + } + + PhOnlineJitFreeCudaSource(cuda_source); + } + result = PhOnlineJitCompileTable(context, table, backend, diff --git a/examples/cpp-console-online-rawdog-jit/CMakeLists.txt b/examples/cpp-console-online-rawdog-jit/CMakeLists.txt index 41cde811..9c839726 100644 --- a/examples/cpp-console-online-rawdog-jit/CMakeLists.txt +++ b/examples/cpp-console-online-rawdog-jit/CMakeLists.txt @@ -45,6 +45,10 @@ set( set(PH_ONLINE_RAWDOG_TARGET "") set(PH_ONLINE_RAWDOG_RUNTIME_TARGET "") +if(PERFECTHASH_ROOT AND NOT PERFECTHASH_ROOT STREQUAL "") + set(PH_ONLINE_RAWDOG_USE_FETCHCONTENT OFF CACHE BOOL "" FORCE) +endif() + if(PH_ONLINE_RAWDOG_USE_FETCHCONTENT) include(FetchContent) diff --git a/examples/cpp-console-online-rawdog-jit/cmake/FindPerfectHashOnlineRawdog.cmake b/examples/cpp-console-online-rawdog-jit/cmake/FindPerfectHashOnlineRawdog.cmake index 27f63a34..0a800827 100644 --- a/examples/cpp-console-online-rawdog-jit/cmake/FindPerfectHashOnlineRawdog.cmake +++ b/examples/cpp-console-online-rawdog-jit/cmake/FindPerfectHashOnlineRawdog.cmake @@ -29,10 +29,13 @@ list(REMOVE_DUPLICATES _ph_search_hints) find_path( PERFECTHASH_ONLINE_RAWDOG_INCLUDE_DIR - NAMES PerfectHashOnlineRawdog.h + NAMES + PerfectHash/PerfectHashOnlineRawdog.h + PerfectHashOnlineRawdog.h HINTS ${_ph_search_hints} PATH_SUFFIXES include + . ) set(_ph_lib_suffixes) diff --git a/examples/tpch-query-probes/README.md b/examples/tpch-query-probes/README.md new file mode 100644 index 00000000..13e40534 --- /dev/null +++ b/examples/tpch-query-probes/README.md @@ -0,0 +1,28 @@ +# tpch-query-probes + +Query-driven extractor for TPC-H SF10 parquet data that materializes: + +1. a unique build-side `.keys` domain for a candidate join, +2. a duplicate-preserving probe stream file for the corresponding fact-side key + accesses, and +3. summary JSON with probe cardinality and skew statistics. + +Current query/join targets: + +- `q5_supplier_lineitem` +- `q8_part_lineitem` +- `q21_supplier_lineitem` + +The output is intended for downstream PerfectHash and cuCollections lookup +benchmarks on realistic probe streams instead of the current unique-key +microbenchmarks. + +## Usage + +```bash +PYTHONPATH=/tmp/tpch_duckdb python examples/tpch-query-probes/extract_tpch_query_probes.py \ + --dataset-root /mnt/data/tpch/scale-10 \ + --output-root /mnt/data/tpch/scale-10/query-probes +``` + +If `duckdb` is installed elsewhere, the script will use the normal import path. diff --git a/examples/tpch-query-probes/extract_tpch_query_probes.py b/examples/tpch-query-probes/extract_tpch_query_probes.py new file mode 100644 index 00000000..f7f1d166 --- /dev/null +++ b/examples/tpch-query-probes/extract_tpch_query_probes.py @@ -0,0 +1,443 @@ +#!/usr/bin/env python3 + +from __future__ import annotations + +import argparse +import csv +import json +import os +import struct +import sys +from pathlib import Path + + +def import_duckdb(): + try: + import duckdb # type: ignore + + return duckdb + except ModuleNotFoundError: + fallback = "/tmp/tpch_duckdb" + if os.path.isdir(fallback): + sys.path.insert(0, fallback) + import duckdb # type: ignore + + return duckdb + raise + + +def write_u64_file(path: Path, values) -> int: + count = 0 + with path.open("wb") as f: + for value in values: + f.write(struct.pack(" dict: + rows = con.execute( + f""" + select distinct k + from ( + {sql} + ) + order by k + """ + ).fetchall() + values = [row[0] for row in rows] + count = write_u64_file(output_path, values) + return { + "path": str(output_path), + "count": count, + "bytes": count * 8, + } + + +def materialize_probe_stream(con, output_path: Path, sql: str) -> dict: + count = 0 + with output_path.open("wb") as f: + cur = con.execute(sql) + while True: + rows = cur.fetchmany(100_000) + if not rows: + break + for (value,) in rows: + f.write(struct.pack(" dict: + summary_sql = f""" + with freq as ( + {sql} + ), ranked as ( + select + k, + cnt, + row_number() over(order by cnt desc, k) as rn, + count(*) over() as dk, + sum(cnt) over(order by cnt desc, k rows between unbounded preceding and current row) as prefix, + sum(cnt) over() as total_rows, + max(cnt) over() as max_cnt + from freq + ) + select + max(total_rows) as total_rows, + max(dk) as distinct_keys, + max(max_cnt) as max_cnt, + max(prefix) filter (where rn <= ceil(dk * 0.01)) as top1_rows, + max(prefix) filter (where rn <= ceil(dk * 0.10)) as top10_rows + from ranked + """ + total_rows, distinct_keys, max_cnt, top1_rows, top10_rows = con.execute(summary_sql).fetchone() + total_rows = int(total_rows or 0) + distinct_keys = int(distinct_keys or 0) + max_cnt = int(max_cnt or 0) + top1_rows = int(top1_rows or 0) + top10_rows = int(top10_rows or 0) + return { + "probe_rows": total_rows, + "probe_distinct_keys": distinct_keys, + "max_frequency": max_cnt, + "avg_multiplicity": (float(total_rows) / float(distinct_keys)) if distinct_keys else 0.0, + "top1_mass": (float(top1_rows) / float(total_rows)) if total_rows else 0.0, + "top10_mass": (float(top10_rows) / float(total_rows)) if total_rows else 0.0, + } + + +def write_top_distinct_subset(con, output_path: Path, sql: str, fraction: float) -> dict: + fraction_sql = f"{fraction:.12f}" + subset_sql = f""" + with freq as ( + {sql} + ), ranked as ( + select + k, + cnt, + row_number() over(order by cnt desc, k) as rn, + count(*) over() as dk, + greatest(1, cast(ceil(count(*) over() * {fraction_sql}) as bigint)) as selected_count + from freq + ) + select k + from ranked + where rn <= selected_count + order by k + """ + rows = con.execute(subset_sql).fetchall() + values = [row[0] for row in rows] + count = write_u64_file(output_path, values) + return { + "path": str(output_path), + "count": count, + "bytes": count * 8, + "fraction_of_distinct": fraction, + "selection_mode": "ceil_with_minimum_one", + "minimum_distinct_keys": 1, + } + + +def register_views(con, dataset_root: Path, tables): + def quote_sql_path(path: Path) -> str: + return str(path).replace("'", "''") + + for table in tables: + path = dataset_root / table / "*.parquet" + con.execute( + f"create or replace view {table} as " + f"select * from parquet_scan('{quote_sql_path(path)}')" + ) + + +def extract_q8(con, output_root: Path): + name = "q8_part_lineitem" + build_sql = """ + select p_partkey as k + from part + where p_type = 'ECONOMY ANODIZED STEEL' + order by k + """ + probe_stream_sql = """ + with orders_region as ( + select o_orderkey + from orders + join customer on o_custkey = c_custkey + join nation on c_nationkey = n_nationkey + join region on n_regionkey = r_regionkey + where r_name = 'AMERICA' + and o_orderdate between date '1995-01-01' and date '1996-12-31' + ) + select l_partkey as k + from lineitem + join orders_region on l_orderkey = o_orderkey + """ + probe_freq_sql = """ + with orders_region as ( + select o_orderkey + from orders + join customer on o_custkey = c_custkey + join nation on c_nationkey = n_nationkey + join region on n_regionkey = r_regionkey + where r_name = 'AMERICA' + and o_orderdate between date '1995-01-01' and date '1996-12-31' + ) + select l_partkey as k, count(*) as cnt + from lineitem + join orders_region on l_orderkey = o_orderkey + group by l_partkey + """ + + target_dir = output_root / name + target_dir.mkdir(parents=True, exist_ok=True) + + build_info = materialize_build_keys(con, target_dir / "build_p_partkey-64.keys", build_sql) + probe_info = materialize_probe_stream(con, target_dir / "probe_l_partkey-64.bin", probe_stream_sql) + summary = summarize_probe_freq(con, probe_freq_sql) + summary["candidate"] = name + summary["build_rows"] = build_info["count"] + summary["probe_to_build_ratio"] = ( + summary["probe_rows"] / build_info["count"] if build_info["count"] else 0.0 + ) + summary["build_keys"] = build_info + summary["probe_stream"] = probe_info + summary["top1pct_distinct_keys"] = write_top_distinct_subset( + con, target_dir / "top1pct_distinct_l_partkey-64.keys", probe_freq_sql, 0.01 + ) + summary["top10pct_distinct_keys"] = write_top_distinct_subset( + con, target_dir / "top10pct_distinct_l_partkey-64.keys", probe_freq_sql, 0.10 + ) + with (target_dir / "summary.json").open("w") as f: + json.dump(summary, f, indent=2, sort_keys=True) + return summary + + +def extract_q21(con, output_root: Path): + name = "q21_supplier_lineitem" + build_sql = """ + select s_suppkey as k + from supplier + join nation on s_nationkey = n_nationkey + where n_name = 'SAUDI ARABIA' + order by k + """ + probe_stream_sql = """ + with multi_supplier_orders as ( + select l_orderkey + from lineitem + group by l_orderkey + having count(l_suppkey) > 1 + ), late_lineitems as ( + select li.l_orderkey, li.l_suppkey + from lineitem li + join multi_supplier_orders mso on li.l_orderkey = mso.l_orderkey + where li.l_receiptdate > li.l_commitdate + ), late_supplier_counts as ( + select l_orderkey, count(l_suppkey) as n_supp_by_order + from late_lineitems + group by l_orderkey + ), saudi_suppliers as ( + select s_suppkey + from supplier + join nation on s_nationkey = n_nationkey + where n_name = 'SAUDI ARABIA' + ) + select lli.l_suppkey as k + from late_supplier_counts lsc + join late_lineitems lli on lsc.l_orderkey = lli.l_orderkey + join orders o on lli.l_orderkey = o.o_orderkey + join saudi_suppliers ss on lli.l_suppkey = ss.s_suppkey + where lsc.n_supp_by_order = 1 + and o.o_orderstatus = 'F' + """ + probe_freq_sql = """ + with multi_supplier_orders as ( + select l_orderkey + from lineitem + group by l_orderkey + having count(l_suppkey) > 1 + ), late_lineitems as ( + select li.l_orderkey, li.l_suppkey + from lineitem li + join multi_supplier_orders mso on li.l_orderkey = mso.l_orderkey + where li.l_receiptdate > li.l_commitdate + ), late_supplier_counts as ( + select l_orderkey, count(l_suppkey) as n_supp_by_order + from late_lineitems + group by l_orderkey + ), saudi_suppliers as ( + select s_suppkey + from supplier + join nation on s_nationkey = n_nationkey + where n_name = 'SAUDI ARABIA' + ) + select lli.l_suppkey as k, count(*) as cnt + from late_supplier_counts lsc + join late_lineitems lli on lsc.l_orderkey = lli.l_orderkey + join orders o on lli.l_orderkey = o.o_orderkey + join saudi_suppliers ss on lli.l_suppkey = ss.s_suppkey + where lsc.n_supp_by_order = 1 + and o.o_orderstatus = 'F' + group by lli.l_suppkey + """ + + target_dir = output_root / name + target_dir.mkdir(parents=True, exist_ok=True) + + build_info = materialize_build_keys(con, target_dir / "build_s_suppkey-64.keys", build_sql) + probe_info = materialize_probe_stream(con, target_dir / "probe_l_suppkey-64.bin", probe_stream_sql) + summary = summarize_probe_freq(con, probe_freq_sql) + summary["candidate"] = name + summary["build_rows"] = build_info["count"] + summary["probe_to_build_ratio"] = ( + summary["probe_rows"] / build_info["count"] if build_info["count"] else 0.0 + ) + summary["build_keys"] = build_info + summary["probe_stream"] = probe_info + summary["top1pct_distinct_keys"] = write_top_distinct_subset( + con, target_dir / "top1pct_distinct_l_suppkey-64.keys", probe_freq_sql, 0.01 + ) + summary["top10pct_distinct_keys"] = write_top_distinct_subset( + con, target_dir / "top10pct_distinct_l_suppkey-64.keys", probe_freq_sql, 0.10 + ) + with (target_dir / "summary.json").open("w") as f: + json.dump(summary, f, indent=2, sort_keys=True) + return summary + + +def extract_q5(con, output_root: Path): + name = "q5_supplier_lineitem" + build_sql = """ + select s_suppkey as k + from supplier + join nation on s_nationkey = n_nationkey + join region on n_regionkey = r_regionkey + where r_name = 'ASIA' + order by k + """ + probe_stream_sql = """ + with asia_customers as ( + select c_custkey + from customer + join nation on c_nationkey = n_nationkey + join region on n_regionkey = r_regionkey + where r_name = 'ASIA' + ), asia_orders as ( + select o_orderkey + from orders + join asia_customers on o_custkey = c_custkey + where o_orderdate >= date '1994-01-01' + and o_orderdate < date '1995-01-01' + ) + select l_suppkey as k + from lineitem + join asia_orders on l_orderkey = o_orderkey + """ + probe_freq_sql = """ + with asia_customers as ( + select c_custkey + from customer + join nation on c_nationkey = n_nationkey + join region on n_regionkey = r_regionkey + where r_name = 'ASIA' + ), asia_orders as ( + select o_orderkey + from orders + join asia_customers on o_custkey = c_custkey + where o_orderdate >= date '1994-01-01' + and o_orderdate < date '1995-01-01' + ) + select l_suppkey as k, count(*) as cnt + from lineitem + join asia_orders on l_orderkey = o_orderkey + group by l_suppkey + """ + + target_dir = output_root / name + target_dir.mkdir(parents=True, exist_ok=True) + + build_info = materialize_build_keys(con, target_dir / "build_s_suppkey-64.keys", build_sql) + probe_info = materialize_probe_stream(con, target_dir / "probe_l_suppkey-64.bin", probe_stream_sql) + summary = summarize_probe_freq(con, probe_freq_sql) + summary["candidate"] = name + summary["build_rows"] = build_info["count"] + summary["probe_to_build_ratio"] = ( + summary["probe_rows"] / build_info["count"] if build_info["count"] else 0.0 + ) + summary["build_keys"] = build_info + summary["probe_stream"] = probe_info + summary["top1pct_distinct_keys"] = write_top_distinct_subset( + con, target_dir / "top1pct_distinct_l_suppkey-64.keys", probe_freq_sql, 0.01 + ) + summary["top10pct_distinct_keys"] = write_top_distinct_subset( + con, target_dir / "top10pct_distinct_l_suppkey-64.keys", probe_freq_sql, 0.10 + ) + with (target_dir / "summary.json").open("w") as f: + json.dump(summary, f, indent=2, sort_keys=True) + return summary + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument("--dataset-root", required=True) + parser.add_argument("--output-root", required=True) + args = parser.parse_args() + + duckdb = import_duckdb() + con = duckdb.connect() + register_views( + con, + Path(args.dataset_root), + ["customer", "orders", "lineitem", "part", "supplier", "nation", "region"], + ) + + output_root = Path(args.output_root) + output_root.mkdir(parents=True, exist_ok=True) + + results = [ + extract_q5(con, output_root), + extract_q8(con, output_root), + extract_q21(con, output_root), + ] + + table_path = output_root / "summary.csv" + with table_path.open("w", newline="") as f: + writer = csv.DictWriter( + f, + fieldnames=[ + "candidate", + "build_rows", + "probe_rows", + "probe_distinct_keys", + "max_frequency", + "avg_multiplicity", + "top1_mass", + "top10_mass", + "probe_to_build_ratio", + ], + ) + writer.writeheader() + for row in results: + writer.writerow({k: row[k] for k in writer.fieldnames}) + + print(table_path) + for row in results: + print( + row["candidate"], + row["build_rows"], + row["probe_rows"], + row["probe_distinct_keys"], + f"avg_mult={row['avg_multiplicity']:.3f}", + f"top1={row['top1_mass']:.3f}", + f"top10={row['top10_mass']:.3f}", + f"probe/build={row['probe_to_build_ratio']:.3f}", + ) + + +if __name__ == "__main__": + main() diff --git a/examples/tpch-query-probes/partition_hot_subset.py b/examples/tpch-query-probes/partition_hot_subset.py new file mode 100644 index 00000000..cb819d8e --- /dev/null +++ b/examples/tpch-query-probes/partition_hot_subset.py @@ -0,0 +1,184 @@ +#!/usr/bin/env python3 + +from __future__ import annotations + +import argparse +import json +import math +import struct +from pathlib import Path + + +def read_u64_file(path: Path): + with path.open("rb") as f: + data = f.read() + if len(data) % 8 != 0: + raise ValueError(f"{path} size is not a multiple of 8") + return [value[0] for value in struct.iter_unpack(" dict[int, int]: + counts = {key: 0 for key in build_keys} + with probe_path.open("rb") as src: + while True: + chunk = src.read(8 * 100_000) + if not chunk: + break + for (value,) in struct.iter_unpack(" Path: + path = Path(path_text) + if path.is_absolute(): + return path + return candidate_dir / path + + build_keys = set(read_u64_file(resolve_summary_path(summary["build_keys"]["path"]))) + probe_count = summary["probe_stream"]["count"] + probe_path = resolve_summary_path(summary["probe_stream"]["path"]) + + for subset_name in ["top1pct_distinct_keys", "top10pct_distinct_keys"]: + subset_info = summary[subset_name] + probe_hot_distinct = read_u64_file(resolve_summary_path(subset_info["path"])) + build_hot_distinct = sorted(set(probe_hot_distinct).intersection(build_keys)) + + tag = "top10pct" if subset_name.startswith("top10") else "top1pct" + build_hot_path = candidate_dir / f"build_hot_{tag}-64.keys" + hot_probe_path = candidate_dir / f"probe_hot_{tag}-64.bin" + cold_probe_path = candidate_dir / f"probe_cold_{tag}-64.bin" + + build_hot_count = write_u64_file(build_hot_path, build_hot_distinct) + hot_probe_count, cold_probe_count = partition_probe_stream( + probe_path, set(build_hot_distinct), hot_probe_path, cold_probe_path + ) + + derived = { + "tag": tag, + "build_hot_keys": { + "path": str(build_hot_path), + "count": build_hot_count, + "bytes": build_hot_count * 8, + }, + "probe_hot_stream": { + "path": str(hot_probe_path), + "count": hot_probe_count, + "bytes": hot_probe_count * 8, + }, + "probe_cold_stream": { + "path": str(cold_probe_path), + "count": cold_probe_count, + "bytes": cold_probe_count * 8, + }, + "hit_mass": (hot_probe_count / probe_count) if probe_count else 0.0, + "miss_mass": (cold_probe_count / probe_count) if probe_count else 0.0, + "build_hot_fraction_of_build": (build_hot_count / summary["build_rows"]) + if summary["build_rows"] else 0.0, + } + + out_path = candidate_dir / f"{tag}_partition_summary.json" + with out_path.open("w") as f: + json.dump(derived, f, indent=2, sort_keys=True) + + print(candidate_dir.name, tag, build_hot_count, hot_probe_count, cold_probe_count, + f"hit_mass={derived['hit_mass']:.4f}", + f"build_hot_frac={derived['build_hot_fraction_of_build']:.4f}") + + counts = summarize_build_key_frequencies(probe_path, build_keys) + ranked_build_hits = sorted(counts.items(), key=lambda kv: (-kv[1], kv[0])) + build_count = summary["build_rows"] + probe_count = summary["probe_stream"]["count"] + + for fraction, tag in [(0.01, "buildhit_top1pct"), (0.10, "buildhit_top10pct")]: + keep = max(1, math.ceil(build_count * fraction)) + hot_keys = sorted(key for key, _ in ranked_build_hits[:keep] if _ > 0) + + build_hot_path = candidate_dir / f"build_hot_{tag}-64.keys" + hot_probe_path = candidate_dir / f"probe_hot_{tag}-64.bin" + cold_probe_path = candidate_dir / f"probe_cold_{tag}-64.bin" + + build_hot_count = write_u64_file(build_hot_path, hot_keys) + hot_probe_count, cold_probe_count = partition_probe_stream( + probe_path, set(hot_keys), hot_probe_path, cold_probe_path + ) + + derived = { + "tag": tag, + "fraction_of_build": fraction, + "build_hot_keys": { + "path": str(build_hot_path), + "count": build_hot_count, + "bytes": build_hot_count * 8, + }, + "probe_hot_stream": { + "path": str(hot_probe_path), + "count": hot_probe_count, + "bytes": hot_probe_count * 8, + }, + "probe_cold_stream": { + "path": str(cold_probe_path), + "count": cold_probe_count, + "bytes": cold_probe_count * 8, + }, + "hit_mass": (hot_probe_count / probe_count) if probe_count else 0.0, + "miss_mass": (cold_probe_count / probe_count) if probe_count else 0.0, + "build_hot_fraction_of_build": (build_hot_count / build_count) if build_count else 0.0, + "min_hot_frequency": min((counts[k] for k in hot_keys), default=0), + "max_hot_frequency": max((counts[k] for k in hot_keys), default=0), + } + + out_path = candidate_dir / f"{tag}_partition_summary.json" + with out_path.open("w") as f: + json.dump(derived, f, indent=2, sort_keys=True) + + print(candidate_dir.name, tag, build_hot_count, hot_probe_count, cold_probe_count, + f"hit_mass={derived['hit_mass']:.4f}", + f"build_hot_frac={derived['build_hot_fraction_of_build']:.4f}", + f"min_hot_freq={derived['min_hot_frequency']}", + f"max_hot_freq={derived['max_hot_frequency']}") + + +if __name__ == "__main__": + main() diff --git a/include/PerfectHash/PerfectHashOnlineJit.h b/include/PerfectHash/PerfectHashOnlineJit.h index 7e9a62ea..3694b1fb 100644 --- a/include/PerfectHash/PerfectHashOnlineJit.h +++ b/include/PerfectHash/PerfectHashOnlineJit.h @@ -44,6 +44,24 @@ extern "C" { typedef struct PH_ONLINE_JIT_CONTEXT PH_ONLINE_JIT_CONTEXT; typedef struct PH_ONLINE_JIT_TABLE PH_ONLINE_JIT_TABLE; +typedef struct PH_ONLINE_JIT_TABLE_INFO { + uint32_t KeySizeInBytes; + uint32_t OriginalKeySizeInBytes; + uint32_t AssignedElementSizeInBytes; + uint32_t HashFunctionId; + uint32_t MaskFunctionId; + uint32_t HashMask; + uint32_t IndexMask; + uint32_t Seed1; + uint32_t Seed2; + uint32_t Seed3; + uint32_t Seed4; + uint32_t Seed5; + uint64_t NumberOfKeys; + uint64_t NumberOfTableElements; + uint64_t DownsizeBitmap; +} PH_ONLINE_JIT_TABLE_INFO; + typedef enum PH_ONLINE_JIT_HASH_FUNCTION { PhOnlineJitHashMultiplyShiftR = 0, PhOnlineJitHashMultiplyShiftLR, @@ -73,6 +91,8 @@ typedef enum PH_ONLINE_JIT_MAX_ISA { } PH_ONLINE_JIT_MAX_ISA; #define PH_ONLINE_JIT_COMPILE_FLAG_STRICT_VECTOR_WIDTH (1u << 0) +#define PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_KERNELS (1u << 0) +#define PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA (1u << 1) PH_ONLINE_JIT_API int32_t @@ -129,6 +149,52 @@ PhOnlineJitCompileTableEx( uint32_t *EffectiveVectorWidth ); +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetCudaSource( + PH_ONLINE_JIT_TABLE *Table, + char **SourceText, + size_t *SourceSize + ); + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetCudaSourceEx( + PH_ONLINE_JIT_TABLE *Table, + uint32_t Flags, + char **SourceText, + size_t *SourceSize + ); + +PH_ONLINE_JIT_API +void +PhOnlineJitFreeCudaSource( + char *SourceText + ); + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetCudaTableData( + PH_ONLINE_JIT_TABLE *Table, + void **TableData, + size_t *TableDataSize, + uint32_t *TableDataElementSize, + size_t *NumberOfTableElements + ); + +PH_ONLINE_JIT_API +void +PhOnlineJitFreeCudaTableData( + void *TableData + ); + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetTableInfo( + PH_ONLINE_JIT_TABLE *Table, + PH_ONLINE_JIT_TABLE_INFO *TableInfo + ); + PH_ONLINE_JIT_API int32_t PhOnlineJitIndex32( @@ -145,6 +211,14 @@ PhOnlineJitIndex64( uint32_t *Index ); +PH_ONLINE_JIT_API +int32_t +PhOnlineJitIndex32x2( + PH_ONLINE_JIT_TABLE *Table, + const uint32_t *Keys, + uint32_t *Indexes + ); + PH_ONLINE_JIT_API int32_t PhOnlineJitIndex32x4( diff --git a/src/PerfectHash/Chm01FileWorkCudaSourceFile.c b/src/PerfectHash/Chm01FileWorkCudaSourceFile.c index 6456371c..e3ad5b6e 100644 --- a/src/PerfectHash/Chm01FileWorkCudaSourceFile.c +++ b/src/PerfectHash/Chm01FileWorkCudaSourceFile.c @@ -59,13 +59,12 @@ PrepareCudaSourceFileChm01( // BaseSize = Context->SystemAllocationGranularity; - TableElementBytes = 16; - KeyElementBytes = ( - Keys->OriginalKeySizeType == LongLongType ? 24 : 16 - ); + TableElementBytes = 64; + KeyElementBytes = 40; RequiredSize = ( BaseSize + + (1024 * 1024) + (NumberOfTableElements * TableElementBytes) + (NumberOfKeys * KeyElementBytes) ); @@ -100,7 +99,6 @@ SaveCudaSourceFileChm01( PULONG Long; ULONGLONG Index; PCSTRING Name; - PULONG Seeds; ULONG Seed1; ULONG Seed2; ULONG Seed3; @@ -110,7 +108,6 @@ SaveCudaSourceFileChm01( ULONG Seed3Byte2; ULONG Seed3Byte3; ULONG Seed3Byte4; - PGRAPH Graph; PULONG Source; PUSHORT Source16; ULONGLONG NumberOfKeys; @@ -122,6 +119,7 @@ SaveCudaSourceFileChm01( PPERFECT_HASH_PATH Path; PPERFECT_HASH_FILE File; PPERFECT_HASH_TABLE Table; + PGRAPH Graph; PTABLE_INFO_ON_DISK TableInfo; ULONGLONG TotalNumberOfElements; const ULONG Indent = 0x20202020; @@ -137,12 +135,11 @@ SaveCudaSourceFileChm01( Path = GetActivePath(File); Name = &Path->TableNameA; TableInfo = Table->TableInfoOnDisk; - TotalNumberOfElements = TableInfo->NumberOfTableElements.QuadPart; Graph = (PGRAPH)Context->SolvedContext; - NumberOfSeeds = Graph->NumberOfSeeds; - Seeds = &Graph->FirstSeed; + TotalNumberOfElements = TableInfo->NumberOfTableElements.QuadPart; + NumberOfSeeds = TableInfo->NumberOfSeeds; NumberOfKeys = Keys->NumberOfKeys.QuadPart; - UsingAssigned16 = IsUsingAssigned16(Graph); + UsingAssigned16 = (TableInfo->AssignedElementSizeInBytes == sizeof(USHORT)); Supported = ( Table->MaskFunctionId == PerfectHashAndMaskFunctionId && IsGoodPerfectHashHashFunctionId(Table->HashFunctionId) @@ -156,12 +153,17 @@ SaveCudaSourceFileChm01( // OUTPUT_RAW("// Auto-generated by Perfect Hash.\n"); + OUTPUT_RAW("// Define PERFECTHASH_GENERATED_CUDA_LIBRARY_ONLY to omit the\n"); + OUTPUT_RAW("// host-side smoke-test harness and emit only reusable lookup\n"); + OUTPUT_RAW("// code suitable for downstream compilation or runtime JIT.\n"); OUTPUT_RAW("#include \n"); OUTPUT_RAW("#include \n"); OUTPUT_RAW("#include \n"); + OUTPUT_RAW("#ifndef PERFECTHASH_GENERATED_CUDA_LIBRARY_ONLY\n"); OUTPUT_RAW("#include \n"); OUTPUT_RAW("#include \n"); - OUTPUT_RAW("#include \n\n"); + OUTPUT_RAW("#include \n"); + OUTPUT_RAW("#endif\n\n"); OUTPUT_RAW("namespace perfecthash::generated::"); OUTPUT_STRING(Name); @@ -252,26 +254,26 @@ SaveCudaSourceFileChm01( OUTPUT_INT(NumberOfSeeds); OUTPUT_RAW("u;\n"); OUTPUT_RAW("inline constexpr std::uint32_t seed1 = "); - OUTPUT_HEX((NumberOfSeeds >= 1) ? Graph->Seeds[0] : 0); + OUTPUT_HEX((NumberOfSeeds >= 1) ? TableInfo->Seed1 : 0); OUTPUT_RAW("u;\n"); OUTPUT_RAW("inline constexpr std::uint32_t seed2 = "); - OUTPUT_HEX((NumberOfSeeds >= 2) ? Graph->Seeds[1] : 0); + OUTPUT_HEX((NumberOfSeeds >= 2) ? TableInfo->Seed2 : 0); OUTPUT_RAW("u;\n"); OUTPUT_RAW("inline constexpr std::uint32_t seed3 = "); - OUTPUT_HEX((NumberOfSeeds >= 3) ? Graph->Seeds[2] : 0); + OUTPUT_HEX((NumberOfSeeds >= 3) ? TableInfo->Seed3 : 0); OUTPUT_RAW("u;\n\n"); OUTPUT_RAW("inline constexpr std::uint32_t seed4 = "); - OUTPUT_HEX((NumberOfSeeds >= 4) ? Graph->Seeds[3] : 0); + OUTPUT_HEX((NumberOfSeeds >= 4) ? TableInfo->Seed4 : 0); OUTPUT_RAW("u;\n"); OUTPUT_RAW("inline constexpr std::uint32_t seed5 = "); - OUTPUT_HEX((NumberOfSeeds >= 5) ? Graph->Seeds[4] : 0); + OUTPUT_HEX((NumberOfSeeds >= 5) ? TableInfo->Seed5 : 0); OUTPUT_RAW("u;\n\n"); - Seed1 = (NumberOfSeeds >= 1 ? Graph->Seeds[0] : 0); - Seed2 = (NumberOfSeeds >= 2 ? Graph->Seeds[1] : 0); - Seed3 = (NumberOfSeeds >= 3 ? Graph->Seeds[2] : 0); - Seed4 = (NumberOfSeeds >= 4 ? Graph->Seeds[3] : 0); - Seed5 = (NumberOfSeeds >= 5 ? Graph->Seeds[4] : 0); + Seed1 = (NumberOfSeeds >= 1 ? TableInfo->Seed1 : 0); + Seed2 = (NumberOfSeeds >= 2 ? TableInfo->Seed2 : 0); + Seed3 = (NumberOfSeeds >= 3 ? TableInfo->Seed3 : 0); + Seed4 = (NumberOfSeeds >= 4 ? TableInfo->Seed4 : 0); + Seed5 = (NumberOfSeeds >= 5 ? TableInfo->Seed5 : 0); Seed3Byte1 = (Seed3 & 0xff); Seed3Byte2 = ((Seed3 >> 8) & 0xff); Seed3Byte3 = ((Seed3 >> 16) & 0xff); @@ -298,7 +300,10 @@ SaveCudaSourceFileChm01( OUTPUT_RAW("number_of_table_elements] = {\n"); if (UsingAssigned16) { - Source16 = Graph->Assigned16; + Source16 = (Graph && Graph->Assigned16) ? Graph->Assigned16 : Table->Assigned16; + if (!ARGUMENT_PRESENT(Source16)) { + return E_POINTER; + } for (Index = 0, Count = 0; Index < TotalNumberOfElements; @@ -321,7 +326,10 @@ SaveCudaSourceFileChm01( } } else { - Source = Graph->Assigned; + Source = (Graph && Graph->Assigned) ? Graph->Assigned : Table->Assigned; + if (!ARGUMENT_PRESENT(Source)) { + return E_POINTER; + } for (Index = 0, Count = 0; Index < TotalNumberOfElements; @@ -675,6 +683,8 @@ SaveCudaSourceFileChm01( // Host-side test entry point. // + OUTPUT_RAW("#ifndef PERFECTHASH_GENERATED_CUDA_LIBRARY_ONLY\n\n"); + OUTPUT_RAW("namespace {\n\n"); OUTPUT_RAW("inline void check_cuda(cudaError_t err, const char* what) {\n"); OUTPUT_RAW(" if (err != cudaSuccess) {\n"); @@ -775,6 +785,7 @@ SaveCudaSourceFileChm01( OUTPUT_RAW(" std::printf(\"ok\\n\");\n"); OUTPUT_RAW(" return 0;\n"); OUTPUT_RAW("}\n"); + OUTPUT_RAW("\n#endif // PERFECTHASH_GENERATED_CUDA_LIBRARY_ONLY\n"); // // Update the number of bytes written. diff --git a/src/PerfectHash/PerfectHash.def b/src/PerfectHash/PerfectHash.def index b0825caa..1689f051 100644 --- a/src/PerfectHash/PerfectHash.def +++ b/src/PerfectHash/PerfectHash.def @@ -7,3 +7,10 @@ EXPORTS PerfectHashPrintError PerfectHashPrintCuError PerfectHashPrintMessage + PhOnlineJitGetCudaSource + PhOnlineJitGetCudaSourceEx + PhOnlineJitFreeCudaSource + PhOnlineJitGetCudaTableData + PhOnlineJitFreeCudaTableData + PhOnlineJitGetTableInfo + PhOnlineJitIndex32x2 diff --git a/src/PerfectHash/PerfectHashOnlineJit.c b/src/PerfectHash/PerfectHashOnlineJit.c index 2fc0cb87..44c0e560 100644 --- a/src/PerfectHash/PerfectHashOnlineJit.c +++ b/src/PerfectHash/PerfectHashOnlineJit.c @@ -125,6 +125,42 @@ PhMapOnlineJitHashFunction( return S_OK; } +static +HRESULT +PhMapPerfectHashHashFunctionToOnlineJit( + _In_ PERFECT_HASH_HASH_FUNCTION_ID HashFunctionId, + _Out_ PH_ONLINE_JIT_HASH_FUNCTION *HashFunction + ) +{ + if (!ARGUMENT_PRESENT(HashFunction)) { + return E_POINTER; + } + + if (HashFunctionId == PerfectHashHashMultiplyShiftRFunctionId) { + *HashFunction = PhOnlineJitHashMultiplyShiftR; + } else if (HashFunctionId == PerfectHashHashMultiplyShiftLRFunctionId) { + *HashFunction = PhOnlineJitHashMultiplyShiftLR; + } else if (HashFunctionId == PerfectHashHashMultiplyShiftRMultiplyFunctionId) { + *HashFunction = PhOnlineJitHashMultiplyShiftRMultiply; + } else if (HashFunctionId == PerfectHashHashMultiplyShiftR2FunctionId) { + *HashFunction = PhOnlineJitHashMultiplyShiftR2; + } else if (HashFunctionId == PerfectHashHashMultiplyShiftRXFunctionId) { + *HashFunction = PhOnlineJitHashMultiplyShiftRX; + } else if (HashFunctionId == PerfectHashHashMulshrolate1RXFunctionId) { + *HashFunction = PhOnlineJitHashMulshrolate1RX; + } else if (HashFunctionId == PerfectHashHashMulshrolate2RXFunctionId) { + *HashFunction = PhOnlineJitHashMulshrolate2RX; + } else if (HashFunctionId == PerfectHashHashMulshrolate3RXFunctionId) { + *HashFunction = PhOnlineJitHashMulshrolate3RX; + } else if (HashFunctionId == PerfectHashHashMulshrolate4RXFunctionId) { + *HashFunction = PhOnlineJitHashMulshrolate4RX; + } else { + return E_INVALIDARG; + } + + return S_OK; +} + static HRESULT PhApplyOnlineJitVectorWidth( @@ -216,6 +252,447 @@ PhCompileOnlineJitBackend( &CompileFlags); } +static +HRESULT +PhEstimateCudaSourceBytes( + _In_ PPERFECT_HASH_TABLE Table, + _In_ ULONG Flags, + _Out_ PULONGLONG AllocationSizePointer + ) +{ + PTABLE_INFO_ON_DISK TableInfo; + ULONGLONG NumberOfTableElements; + ULONGLONG TableElementBytes; + ULONGLONG ValuesPerLine; + ULONGLONG CharsPerValue; + ULONGLONG NumberOfLines; + ULONGLONG BaseSize; + ULONGLONG RequiredSize; + + if (!ARGUMENT_PRESENT(Table) || + !ARGUMENT_PRESENT(Table->TableInfoOnDisk) || + !ARGUMENT_PRESENT(AllocationSizePointer)) { + return E_POINTER; + } + + TableInfo = Table->TableInfoOnDisk; + + NumberOfTableElements = TableInfo->NumberOfTableElements.QuadPart; + + BaseSize = (Table->Context && Table->Context->SystemAllocationGranularity) + ? Table->Context->SystemAllocationGranularity + : 65536; + if ((Flags & PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA) != 0) { + TableElementBytes = 0; + } else { + if (TableInfo->AssignedElementSizeInBytes == sizeof(USHORT)) { + ValuesPerLine = 4; + CharsPerValue = 12; + } else { + ValuesPerLine = 4; + CharsPerValue = 12; + } + NumberOfLines = (NumberOfTableElements + ValuesPerLine - 1) / ValuesPerLine; + TableElementBytes = + (NumberOfTableElements * CharsPerValue) + + (NumberOfLines * 4) + + 32; + } + + RequiredSize = ( + BaseSize + + TableElementBytes + + (256 * 1024) + ); + + *AllocationSizePointer = ALIGN_UP(RequiredSize, BaseSize); + + return S_OK; +} + +#define PH_ONLINE_CUDA_INDENT() do { \ + *Output++ = ' '; \ + *Output++ = ' '; \ + *Output++ = ' '; \ + *Output++ = ' '; \ +} while (0) + +static +HRESULT +PhEmitOnlineCudaSource( + _In_ PPERFECT_HASH_TABLE Table, + _In_ PCSTRING Name, + _In_ ULONG Flags, + _Out_ PCHAR Base, + _Out_ PULONGLONG NumberOfBytesWrittenPointer + ) +{ + PCHAR Output; + ULONG Count; + ULONGLONG Index; + ULONG Seed1; + ULONG Seed2; + ULONG Seed3; + ULONG Seed4; + ULONG Seed5; + ULONG Seed3Byte1; + ULONG Seed3Byte2; + ULONG Seed3Byte3; + ULONG Seed3Byte4; + PULONG Source; + PUSHORT Source16; + BOOLEAN UsingAssigned16; + STRING DecimalString; + CHAR NumberOfKeysBuffer[32]; + CHAR NumberOfTableElementsBuffer[32]; + HRESULT Result = S_OK; + PTABLE_INFO_ON_DISK TableInfo; + + if (!ARGUMENT_PRESENT(Table) || + !ARGUMENT_PRESENT(Table->TableInfoOnDisk) || + !ARGUMENT_PRESENT(Name) || + !ARGUMENT_PRESENT(Name->Buffer) || + !ARGUMENT_PRESENT(Base) || + !ARGUMENT_PRESENT(NumberOfBytesWrittenPointer)) { + return E_POINTER; + } + + TableInfo = Table->TableInfoOnDisk; + UsingAssigned16 = (TableInfo->AssignedElementSizeInBytes == sizeof(USHORT)); + Seed1 = (TableInfo->NumberOfSeeds >= 1) ? TableInfo->Seed1 : 0; + Seed2 = (TableInfo->NumberOfSeeds >= 2) ? TableInfo->Seed2 : 0; + Seed3 = (TableInfo->NumberOfSeeds >= 3) ? TableInfo->Seed3 : 0; + Seed4 = (TableInfo->NumberOfSeeds >= 4) ? TableInfo->Seed4 : 0; + Seed5 = (TableInfo->NumberOfSeeds >= 5) ? TableInfo->Seed5 : 0; + Seed3Byte1 = (Seed3 & 0xff); + Seed3Byte2 = ((Seed3 >> 8) & 0xff); + Seed3Byte3 = ((Seed3 >> 16) & 0xff); + Seed3Byte4 = ((Seed3 >> 24) & 0xff); + + Output = Base; + DecimalString.Buffer = NULL; + DecimalString.Length = 0; + DecimalString.MaximumLength = sizeof(NumberOfKeysBuffer); + + OUTPUT_RAW("// Auto-generated by Perfect Hash online JIT.\n"); + OUTPUT_RAW("#ifndef PERFECTHASH_ONLINE_JIT_INDEX_KERNEL_NAME\n"); + OUTPUT_RAW("#define PERFECTHASH_ONLINE_JIT_INDEX_KERNEL_NAME index_kernel\n"); + OUTPUT_RAW("#endif\n\n"); + OUTPUT_RAW("#include \n"); + OUTPUT_RAW("#include \n\n"); + + OUTPUT_RAW("namespace perfecthash::generated::"); + OUTPUT_STRING(Name); + OUTPUT_RAW(" {\n\n"); + + OUTPUT_RAW("inline constexpr uint32_t algorithm_id = "); + OUTPUT_INT(Table->AlgorithmId); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t hash_function_id = "); + OUTPUT_INT(Table->HashFunctionId); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t mask_function_id = "); + OUTPUT_INT(Table->MaskFunctionId); + OUTPUT_RAW("u;\n\n"); + + OUTPUT_RAW("inline constexpr uint32_t key_size_bytes = "); + OUTPUT_INT(TableInfo->KeySizeInBytes); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t original_key_size_bytes = "); + OUTPUT_INT(TableInfo->OriginalKeySizeInBytes); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr size_t number_of_keys = "); + DecimalString.Buffer = NumberOfKeysBuffer; + DecimalString.MaximumLength = sizeof(NumberOfKeysBuffer); + DecimalString.Length = 0; + AppendLongLongIntegerToString(&DecimalString, + TableInfo->NumberOfKeys.QuadPart, + CountNumberOfLongLongDigitsInline( + TableInfo->NumberOfKeys.QuadPart + ), + '\0'); + OUTPUT_STRING(&DecimalString); + OUTPUT_RAW(";\n"); + OUTPUT_RAW("inline constexpr size_t number_of_table_elements = "); + DecimalString.Buffer = NumberOfTableElementsBuffer; + DecimalString.MaximumLength = sizeof(NumberOfTableElementsBuffer); + DecimalString.Length = 0; + AppendLongLongIntegerToString(&DecimalString, + TableInfo->NumberOfTableElements.QuadPart, + CountNumberOfLongLongDigitsInline( + TableInfo->NumberOfTableElements.QuadPart + ), + '\0'); + OUTPUT_STRING(&DecimalString); + OUTPUT_RAW(";\n\n"); + + OUTPUT_RAW("inline constexpr uint32_t hash_mask = "); + OUTPUT_HEX(TableInfo->HashMask); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t index_mask = "); + OUTPUT_HEX(TableInfo->IndexMask); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint64_t downsize_bitmap = "); + OUTPUT_HEX64(Table->DownsizeBitmap); + OUTPUT_RAW("ULL;\n\n"); + + if (TableInfo->KeySizeInBytes <= 4) { + OUTPUT_RAW("using key_type = uint32_t;\n"); + } else { + OUTPUT_RAW("using key_type = uint64_t;\n"); + } + + if (TableInfo->OriginalKeySizeInBytes <= 4) { + OUTPUT_RAW("using original_key_type = uint32_t;\n"); + } else { + OUTPUT_RAW("using original_key_type = uint64_t;\n"); + } + + if (UsingAssigned16) { + OUTPUT_RAW("using table_data_type = uint16_t;\n\n"); + } else { + OUTPUT_RAW("using table_data_type = uint32_t;\n\n"); + } + + if ((Flags & PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA) != 0) { + OUTPUT_RAW("inline constexpr bool has_embedded_table_data = false;\n\n"); + } else { + OUTPUT_RAW("inline constexpr bool has_embedded_table_data = true;\n\n"); + } + + OUTPUT_RAW("inline constexpr uint32_t seed1 = "); + OUTPUT_HEX(Seed1); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed2 = "); + OUTPUT_HEX(Seed2); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed3 = "); + OUTPUT_HEX(Seed3); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed4 = "); + OUTPUT_HEX(Seed4); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed5 = "); + OUTPUT_HEX(Seed5); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed3_byte1 = "); + OUTPUT_INT(Seed3Byte1); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed3_byte2 = "); + OUTPUT_INT(Seed3Byte2); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed3_byte3 = "); + OUTPUT_INT(Seed3Byte3); + OUTPUT_RAW("u;\n"); + OUTPUT_RAW("inline constexpr uint32_t seed3_byte4 = "); + OUTPUT_INT(Seed3Byte4); + OUTPUT_RAW("u;\n\n"); + + if ((Flags & PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA) == 0) { + OUTPUT_RAW("inline constexpr table_data_type table_data[number_of_table_elements] = {\n"); + if (UsingAssigned16) { + Source16 = Table->Assigned16; + for (Index = 0, Count = 0; Index < TableInfo->NumberOfTableElements.QuadPart; Index++) { + if (Count == 0) { + PH_ONLINE_CUDA_INDENT(); + } + OUTPUT_HEX(*Source16++); + *Output++ = ','; + if (++Count == 8) { + Count = 0; + *Output++ = '\n'; + } else { + *Output++ = ' '; + } + } + } else { + Source = Table->Assigned; + for (Index = 0, Count = 0; Index < TableInfo->NumberOfTableElements.QuadPart; Index++) { + if (Count == 0) { + PH_ONLINE_CUDA_INDENT(); + } + OUTPUT_HEX(*Source++); + *Output++ = ','; + if (++Count == 4) { + Count = 0; + *Output++ = '\n'; + } else { + *Output++ = ' '; + } + } + } + if (*(Output - 1) == ' ') { + *(Output - 1) = '\n'; + } + OUTPUT_RAW("};\n\n"); + } else { + OUTPUT_RAW("// table_data is supplied by the consumer at runtime.\n\n"); + } + + OUTPUT_RAW("__host__ __device__ __forceinline__ uint64_t extract_bits64(" + "uint64_t value, uint64_t bitmap) noexcept {\n"); + OUTPUT_RAW(" uint64_t result = 0;\n"); + OUTPUT_RAW(" uint64_t out_bit = 0;\n"); + OUTPUT_RAW(" while (bitmap) {\n"); + OUTPUT_RAW(" const uint64_t lsb = bitmap & (~bitmap + 1);\n"); + OUTPUT_RAW(" if (value & lsb) {\n"); + OUTPUT_RAW(" result |= (1ULL << out_bit);\n"); + OUTPUT_RAW(" }\n"); + OUTPUT_RAW(" bitmap ^= lsb;\n"); + OUTPUT_RAW(" ++out_bit;\n"); + OUTPUT_RAW(" }\n"); + OUTPUT_RAW(" return result;\n"); + OUTPUT_RAW("}\n\n"); + + OUTPUT_RAW("__host__ __device__ __forceinline__ constexpr uint64_t mask_from_bits(uint32_t bits) noexcept {\n"); + OUTPUT_RAW(" return (bits >= 64u) ? ~0ULL : ((1ULL << bits) - 1ULL);\n"); + OUTPUT_RAW("}\n"); + OUTPUT_RAW("inline constexpr uint64_t key_mask = mask_from_bits(key_size_bytes * 8u);\n"); + OUTPUT_RAW("inline constexpr uint64_t original_key_mask = mask_from_bits(original_key_size_bytes * 8u);\n\n"); + + OUTPUT_RAW("__host__ __device__ __forceinline__ key_type downsize_key(" + "uint64_t key) noexcept {\n"); + OUTPUT_RAW(" key &= original_key_mask;\n"); + OUTPUT_RAW(" if (downsize_bitmap) {\n"); + OUTPUT_RAW(" return static_cast(extract_bits64(key, downsize_bitmap) & key_mask);\n"); + OUTPUT_RAW(" }\n"); + OUTPUT_RAW(" return static_cast(key & key_mask);\n"); + OUTPUT_RAW("}\n\n"); + + OUTPUT_RAW("__host__ __device__ __forceinline__ uint32_t rotr32(" + "uint32_t value, uint32_t shift) noexcept {\n"); + OUTPUT_RAW(" shift &= 31u;\n"); + OUTPUT_RAW(" if (shift == 0u) { return value; }\n"); + OUTPUT_RAW(" return (value >> shift) | (value << (32u - shift));\n"); + OUTPUT_RAW("}\n\n"); + + OUTPUT_RAW("struct slot_pair_type {\n"); + OUTPUT_RAW(" uint32_t first;\n"); + OUTPUT_RAW(" uint32_t second;\n"); + OUTPUT_RAW("};\n\n"); + + OUTPUT_RAW("__host__ __device__ __forceinline__ slot_pair_type slot_pair_from_key(" + "original_key_type key) noexcept {\n"); + OUTPUT_RAW(" const key_type downsized = downsize_key(static_cast(key));\n"); + if (Table->HashFunctionId == PerfectHashHashMultiplyShiftRFunctionId) { + OUTPUT_RAW(" key_type vertex1 = static_cast(downsized * static_cast(seed1));\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" key_type vertex2 = static_cast(downsized * static_cast(seed2));\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte2;\n"); + OUTPUT_RAW(" return slot_pair_type{static_cast(vertex1 & hash_mask), static_cast(vertex2 & hash_mask)};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMultiplyShiftLRFunctionId) { + OUTPUT_RAW(" key_type vertex1 = static_cast(downsized * static_cast(seed1));\n"); + OUTPUT_RAW(" vertex1 = static_cast(vertex1 << seed3_byte1);\n"); + OUTPUT_RAW(" key_type vertex2 = static_cast(downsized * static_cast(seed2));\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte2;\n"); + OUTPUT_RAW(" return slot_pair_type{static_cast(vertex1 & hash_mask), static_cast(vertex2 & hash_mask)};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMultiplyShiftRMultiplyFunctionId) { + OUTPUT_RAW(" key_type vertex1 = static_cast(downsized * static_cast(seed1));\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" vertex1 = static_cast(vertex1 * static_cast(seed2));\n"); + OUTPUT_RAW(" key_type vertex2 = static_cast(downsized * static_cast(seed4));\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte2;\n"); + OUTPUT_RAW(" vertex2 = static_cast(vertex2 * static_cast(seed5));\n"); + OUTPUT_RAW(" return slot_pair_type{static_cast(vertex1 & hash_mask), static_cast(vertex2 & hash_mask)};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMultiplyShiftR2FunctionId) { + OUTPUT_RAW(" key_type vertex1 = static_cast(downsized * static_cast(seed1));\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" vertex1 = static_cast(vertex1 * static_cast(seed2));\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte2;\n"); + OUTPUT_RAW(" key_type vertex2 = static_cast(downsized * static_cast(seed4));\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte3;\n"); + OUTPUT_RAW(" vertex2 = static_cast(vertex2 * static_cast(seed5));\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte4;\n"); + OUTPUT_RAW(" return slot_pair_type{static_cast(vertex1 & hash_mask), static_cast(vertex2 & hash_mask)};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMultiplyShiftRXFunctionId) { + OUTPUT_RAW(" key_type vertex1 = static_cast(downsized * static_cast(seed1));\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" key_type vertex2 = static_cast(downsized * static_cast(seed2));\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte1;\n"); + OUTPUT_RAW(" return slot_pair_type{static_cast(vertex1), static_cast(vertex2)};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMulshrolate1RXFunctionId) { + OUTPUT_RAW(" const uint32_t downsized32 = static_cast(downsized);\n"); + OUTPUT_RAW(" uint32_t vertex1 = downsized32 * seed1;\n"); + OUTPUT_RAW(" vertex1 = rotr32(vertex1, seed3_byte2);\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" uint32_t vertex2 = downsized32 * seed2;\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte1;\n"); + OUTPUT_RAW(" return slot_pair_type{vertex1, vertex2};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMulshrolate2RXFunctionId) { + OUTPUT_RAW(" const uint32_t downsized32 = static_cast(downsized);\n"); + OUTPUT_RAW(" uint32_t vertex1 = downsized32 * seed1;\n"); + OUTPUT_RAW(" vertex1 = rotr32(vertex1, seed3_byte2);\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" uint32_t vertex2 = downsized32 * seed2;\n"); + OUTPUT_RAW(" vertex2 = rotr32(vertex2, seed3_byte3);\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte1;\n"); + OUTPUT_RAW(" return slot_pair_type{vertex1, vertex2};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMulshrolate3RXFunctionId) { + OUTPUT_RAW(" const uint32_t downsized32 = static_cast(downsized);\n"); + OUTPUT_RAW(" uint32_t vertex1 = downsized32 * seed1;\n"); + OUTPUT_RAW(" vertex1 = rotr32(vertex1, seed3_byte2);\n"); + OUTPUT_RAW(" vertex1 = vertex1 * seed4;\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" uint32_t vertex2 = downsized32 * seed2;\n"); + OUTPUT_RAW(" vertex2 = rotr32(vertex2, seed3_byte3);\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte1;\n"); + OUTPUT_RAW(" return slot_pair_type{vertex1, vertex2};\n"); + } else if (Table->HashFunctionId == PerfectHashHashMulshrolate4RXFunctionId) { + OUTPUT_RAW(" const uint32_t downsized32 = static_cast(downsized);\n"); + OUTPUT_RAW(" uint32_t vertex1 = downsized32 * seed1;\n"); + OUTPUT_RAW(" vertex1 = rotr32(vertex1, seed3_byte2);\n"); + OUTPUT_RAW(" vertex1 = vertex1 * seed4;\n"); + OUTPUT_RAW(" vertex1 >>= seed3_byte1;\n"); + OUTPUT_RAW(" uint32_t vertex2 = downsized32 * seed2;\n"); + OUTPUT_RAW(" vertex2 = rotr32(vertex2, seed3_byte3);\n"); + OUTPUT_RAW(" vertex2 = vertex2 * seed5;\n"); + OUTPUT_RAW(" vertex2 >>= seed3_byte1;\n"); + OUTPUT_RAW(" return slot_pair_type{vertex1, vertex2};\n"); + } else { + return PH_E_NOT_IMPLEMENTED; + } + OUTPUT_RAW("}\n\n"); + + OUTPUT_RAW("__host__ __device__ __forceinline__ uint32_t index_from_key(" + "original_key_type key, const table_data_type* table) noexcept {\n"); + OUTPUT_RAW(" const slot_pair_type slots = slot_pair_from_key(key);\n"); + OUTPUT_RAW(" const uint32_t value_low = static_cast(table[slots.first]);\n"); + OUTPUT_RAW(" const uint32_t value_high = static_cast(table[slots.second]);\n"); + OUTPUT_RAW(" return static_cast((value_low + value_high) & index_mask);\n"); + OUTPUT_RAW("}\n\n"); + + if ((Flags & PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_KERNELS) == 0) { + if ((Flags & PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA) != 0) { + OUTPUT_RAW("__global__ void PERFECTHASH_ONLINE_JIT_INDEX_KERNEL_NAME(" + "const original_key_type* query_keys, " + "const table_data_type* table_data, " + "uint32_t* out, size_t count) {\n"); + } else { + OUTPUT_RAW("__global__ void PERFECTHASH_ONLINE_JIT_INDEX_KERNEL_NAME(" + "const original_key_type* query_keys, " + "uint32_t* out, size_t count) {\n"); + } + OUTPUT_RAW(" const size_t i = (blockIdx.x * blockDim.x) + threadIdx.x;\n"); + OUTPUT_RAW(" if (i < count) {\n"); + if ((Flags & PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA) != 0) { + OUTPUT_RAW(" out[i] = index_from_key(query_keys[i], table_data);\n"); + } else { + OUTPUT_RAW(" out[i] = index_from_key(query_keys[i], perfecthash::generated::"); + OUTPUT_STRING(Name); + OUTPUT_RAW("::table_data);\n"); + } + OUTPUT_RAW(" }\n"); + OUTPUT_RAW("}\n\n"); + } + + OUTPUT_RAW("} // namespace perfecthash::generated::"); + OUTPUT_STRING(Name); + OUTPUT_RAW("\n"); + + *NumberOfBytesWrittenPointer = RtlPointerToOffset(Base, Output); + return Result; +} + PH_ONLINE_JIT_API int32_t PhOnlineJitOpen( @@ -482,7 +959,7 @@ PhOnlineJitCompileTableEx( { HRESULT Result; HRESULT LastResult; - ULONG CandidateWidths[4] = {0}; + ULONG CandidateWidths[8] = {0}; ULONG CandidateCount = 0; ULONG Index; BOOLEAN StrictVectorWidth; @@ -607,27 +1084,296 @@ PhOnlineJitCompileTableEx( return LastResult; } - Result = PhCompileOnlineJitBackend(Context, - Table, - PhOnlineJitBackendLlvmJit, - VectorWidth, - JitMaxIsa); - if (SUCCEEDED(Result)) { - if (ARGUMENT_PRESENT(EffectiveBackend)) { - *EffectiveBackend = PhOnlineJitBackendLlvmJit; + for (Index = 0; Index < CandidateCount; Index++) { + Result = PhCompileOnlineJitBackend(Context, + Table, + PhOnlineJitBackendLlvmJit, + CandidateWidths[Index], + JitMaxIsa); + if (SUCCEEDED(Result)) { + if (ARGUMENT_PRESENT(EffectiveBackend)) { + *EffectiveBackend = PhOnlineJitBackendLlvmJit; + } + if (ARGUMENT_PRESENT(EffectiveVectorWidth)) { + *EffectiveVectorWidth = CandidateWidths[Index]; + } + return S_OK; } - if (ARGUMENT_PRESENT(EffectiveVectorWidth)) { - *EffectiveVectorWidth = VectorWidth; + + if (Result != PH_E_LLVM_BACKEND_NOT_FOUND && + Result != PH_E_NOT_IMPLEMENTED && + Result != PH_E_INVARIANT_CHECK_FAILED) { + return Result; } - return S_OK; + + LastResult = Result; + } + + return LastResult; +} + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetCudaSource( + PH_ONLINE_JIT_TABLE *Table, + char **SourceText, + size_t *SourceSize + ) +{ + return PhOnlineJitGetCudaSourceEx(Table, 0, SourceText, SourceSize); +} + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetCudaSourceEx( + PH_ONLINE_JIT_TABLE *Table, + uint32_t Flags, + char **SourceText, + size_t *SourceSize + ) +{ + HRESULT Result; + ULONGLONG AllocationSize; + PCHAR Buffer = NULL; + ULONGLONG NumberOfBytesWritten = 0; + STRING Name = {0}; + CHAR NameBuffer[] = "online_jit_table"; + + if (!ARGUMENT_PRESENT(Table) || + !ARGUMENT_PRESENT(Table->Table) || + !ARGUMENT_PRESENT(SourceText)) { + return E_INVALIDARG; + } + if ((Flags & ~(PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_KERNELS | + PH_ONLINE_JIT_CUDA_SOURCE_FLAG_OMIT_TABLE_DATA)) != 0) { + return E_INVALIDARG; + } + + *SourceText = NULL; + if (ARGUMENT_PRESENT(SourceSize)) { + *SourceSize = 0; } - if (Result != PH_E_LLVM_BACKEND_NOT_FOUND && - Result != PH_E_NOT_IMPLEMENTED) { + Result = PhEstimateCudaSourceBytes(Table->Table, Flags, &AllocationSize); + if (FAILED(Result)) { return Result; } + if (AllocationSize > ((ULONGLONG)~((size_t)0)) - 1) { + return E_OUTOFMEMORY; + } - return LastResult; + Buffer = (PCHAR)calloc(1, (size_t)AllocationSize + 1); + if (!Buffer) { + return E_OUTOFMEMORY; + } + + Name.Buffer = NameBuffer; + Name.Length = sizeof(NameBuffer) - sizeof(CHAR); + Name.MaximumLength = sizeof(NameBuffer); + + Result = PhEmitOnlineCudaSource(Table->Table, + &Name, + Flags, + Buffer, + &NumberOfBytesWritten); + if (FAILED(Result)) { + free(Buffer); + return Result; + } + + Buffer[NumberOfBytesWritten] = '\0'; + + *SourceText = Buffer; + if (ARGUMENT_PRESENT(SourceSize)) { + *SourceSize = (size_t)NumberOfBytesWritten; + } + + return S_OK; +} + +PH_ONLINE_JIT_API +void +PhOnlineJitFreeCudaSource( + char *SourceText + ) +{ + free(SourceText); +} + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetCudaTableData( + PH_ONLINE_JIT_TABLE *Table, + void **TableData, + size_t *TableDataSize, + uint32_t *TableDataElementSize, + size_t *NumberOfTableElements + ) +{ + PTABLE_INFO_ON_DISK TableInfo; + ULONGLONG NumberOfTableElements64; + ULONG ElementSizeInBytes; + size_t CopySize; + PVOID Buffer; + PVOID Source; + + if (!ARGUMENT_PRESENT(Table) || + !ARGUMENT_PRESENT(Table->Table) || + !ARGUMENT_PRESENT(TableData) || + !ARGUMENT_PRESENT(Table->Table->TableInfoOnDisk)) { + return E_INVALIDARG; + } + + *TableData = NULL; + if (ARGUMENT_PRESENT(TableDataSize)) { + *TableDataSize = 0; + } + if (ARGUMENT_PRESENT(TableDataElementSize)) { + *TableDataElementSize = 0; + } + if (ARGUMENT_PRESENT(NumberOfTableElements)) { + *NumberOfTableElements = 0; + } + + TableInfo = Table->Table->TableInfoOnDisk; + NumberOfTableElements64 = TableInfo->NumberOfTableElements.QuadPart; + ElementSizeInBytes = TableInfo->AssignedElementSizeInBytes; + if (ElementSizeInBytes == 0) { + return E_INVALIDARG; + } + if (NumberOfTableElements64 > + (((ULONGLONG)~((size_t)0)) / (ULONGLONG)ElementSizeInBytes)) { + return E_OUTOFMEMORY; + } + CopySize = (size_t)(NumberOfTableElements64 * (ULONGLONG)ElementSizeInBytes); + + if (ElementSizeInBytes == sizeof(USHORT)) { + Source = Table->Table->Assigned16; + } else if (ElementSizeInBytes == sizeof(ULONG)) { + Source = Table->Table->Assigned; + } else { + return PH_E_NOT_IMPLEMENTED; + } + + if (Source == NULL) { + return E_POINTER; + } + + Buffer = malloc(CopySize); + if (Buffer == NULL) { + return E_OUTOFMEMORY; + } + + memmove(Buffer, Source, CopySize); + + *TableData = Buffer; + if (ARGUMENT_PRESENT(TableDataSize)) { + *TableDataSize = CopySize; + } + if (ARGUMENT_PRESENT(TableDataElementSize)) { + *TableDataElementSize = ElementSizeInBytes; + } + if (ARGUMENT_PRESENT(NumberOfTableElements)) { + *NumberOfTableElements = (size_t)NumberOfTableElements64; + } + + return S_OK; +} + +PH_ONLINE_JIT_API +void +PhOnlineJitFreeCudaTableData( + void *TableData + ) +{ + free(TableData); +} + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitGetTableInfo( + PH_ONLINE_JIT_TABLE *Table, + PH_ONLINE_JIT_TABLE_INFO *TableInfo + ) +{ + HRESULT Result; + PTABLE_INFO_ON_DISK Info; + + if (!ARGUMENT_PRESENT(Table) || + !ARGUMENT_PRESENT(Table->Table) || + !ARGUMENT_PRESENT(TableInfo) || + !ARGUMENT_PRESENT(Table->Table->TableInfoOnDisk)) { + return E_INVALIDARG; + } + + Info = Table->Table->TableInfoOnDisk; + { + PUCHAR TableInfoBytes = (PUCHAR)TableInfo; + for (size_t Offset = 0; Offset < sizeof(*TableInfo); Offset++) { + TableInfoBytes[Offset] = 0; + } + } + + TableInfo->KeySizeInBytes = Info->KeySizeInBytes; + TableInfo->OriginalKeySizeInBytes = Info->OriginalKeySizeInBytes; + TableInfo->AssignedElementSizeInBytes = Info->AssignedElementSizeInBytes; + { + PH_ONLINE_JIT_HASH_FUNCTION PublicHashFunction; + + Result = PhMapPerfectHashHashFunctionToOnlineJit( + Table->Table->HashFunctionId, + &PublicHashFunction + ); + if (FAILED(Result)) { + return Result; + } + TableInfo->HashFunctionId = (uint32_t)PublicHashFunction; + } + TableInfo->MaskFunctionId = Table->Table->MaskFunctionId; + TableInfo->HashMask = Info->HashMask; + TableInfo->IndexMask = Info->IndexMask; + TableInfo->Seed1 = Info->Seed1; + TableInfo->Seed2 = Info->Seed2; + TableInfo->Seed3 = Info->Seed3; + TableInfo->Seed4 = Info->Seed4; + TableInfo->Seed5 = Info->Seed5; + TableInfo->NumberOfKeys = Info->NumberOfKeys.QuadPart; + TableInfo->NumberOfTableElements = Info->NumberOfTableElements.QuadPart; + TableInfo->DownsizeBitmap = Table->Table->DownsizeBitmap; + + return S_OK; +} + +PH_ONLINE_JIT_API +int32_t +PhOnlineJitIndex32x2( + PH_ONLINE_JIT_TABLE *Table, + const uint32_t *Keys, + uint32_t *Indexes + ) +{ + HRESULT Result; + PPERFECT_HASH_TABLE_JIT_INTERFACE Jit; + + if (!ARGUMENT_PRESENT(Table) || + !ARGUMENT_PRESENT(Keys) || + !ARGUMENT_PRESENT(Indexes)) { + return E_INVALIDARG; + } + + Result = PhEnsureOnlineJitInterface(Table); + if (FAILED(Result)) { + return Result; + } + + Jit = Table->JitInterface; + return Jit->Vtbl->Index32x2( + Jit, + (ULONG)Keys[0], + (ULONG)Keys[1], + (PULONG)&Indexes[0], + (PULONG)&Indexes[1] + ); } PH_ONLINE_JIT_API