From c8c79c1ad9eccbcbae6442f00a18643a7e32448e Mon Sep 17 00:00:00 2001 From: FiveMovesAhead Date: Tue, 2 Dec 2025 16:53:50 +0000 Subject: [PATCH] Submitted vector_search/tree_times_a_lady --- tig-algorithms/src/vector_search/mod.rs | 3 +- .../vector_search/tree_times_a_lady/README.md | 23 ++ .../tree_times_a_lady/kernels.cu | 19 ++ .../vector_search/tree_times_a_lady/mod.rs | 251 ++++++++++++++++++ 4 files changed, 295 insertions(+), 1 deletion(-) create mode 100644 tig-algorithms/src/vector_search/tree_times_a_lady/README.md create mode 100644 tig-algorithms/src/vector_search/tree_times_a_lady/kernels.cu create mode 100644 tig-algorithms/src/vector_search/tree_times_a_lady/mod.rs diff --git a/tig-algorithms/src/vector_search/mod.rs b/tig-algorithms/src/vector_search/mod.rs index ae3472ac..f4836a96 100644 --- a/tig-algorithms/src/vector_search/mod.rs +++ b/tig-algorithms/src/vector_search/mod.rs @@ -42,7 +42,8 @@ // c004_a022 -// c004_a023 +pub mod tree_times_a_lady; +pub use tree_times_a_lady as c004_a023; // c004_a024 diff --git a/tig-algorithms/src/vector_search/tree_times_a_lady/README.md b/tig-algorithms/src/vector_search/tree_times_a_lady/README.md new file mode 100644 index 00000000..7492010e --- /dev/null +++ b/tig-algorithms/src/vector_search/tree_times_a_lady/README.md @@ -0,0 +1,23 @@ +# TIG Code Submission + +## Submission Details + +* **Challenge Name:** vector_search +* **Algorithm Name:** tree_times_a_lady +* **Copyright:** 2024 OvErLoDe +* **Identity of Submitter:** OvErLoDe +* **Identity of Creator of Algorithmic Method:** null +* **Unique Algorithm Identifier (UAI):** null + +## License + +The files in this folder are under the following licenses: +* TIG Benchmarker Outbound License +* TIG Commercial License +* TIG Inbound Game License +* TIG Innovator Outbound Game License +* TIG Open Data License +* TIG THV Game License + +Copies of the licenses can be obtained at: +https://github.com/tig-foundation/tig-monorepo/tree/main/docs/licenses \ No newline at end of file diff --git a/tig-algorithms/src/vector_search/tree_times_a_lady/kernels.cu b/tig-algorithms/src/vector_search/tree_times_a_lady/kernels.cu new file mode 100644 index 00000000..135ce375 --- /dev/null +++ b/tig-algorithms/src/vector_search/tree_times_a_lady/kernels.cu @@ -0,0 +1,19 @@ +/*! +Copyright 2024 OvErLoDe + +Licensed under the TIG Inbound Game License v1.0 or (at your option) any later +version (the "License"); you may not use this file except in compliance with the +License. You may obtain a copy of the License at + +https://github.com/tig-foundation/tig-monorepo/tree/main/docs/licenses + +Unless required by applicable law or agreed to in writing, software distributed +under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR +CONDITIONS OF ANY KIND, either express or implied. See the License for the specific +language governing permissions and limitations under the License. +*/ + +extern "C" __global__ void do_nothing() +{ + // This kernel does nothing +} diff --git a/tig-algorithms/src/vector_search/tree_times_a_lady/mod.rs b/tig-algorithms/src/vector_search/tree_times_a_lady/mod.rs new file mode 100644 index 00000000..cc205982 --- /dev/null +++ b/tig-algorithms/src/vector_search/tree_times_a_lady/mod.rs @@ -0,0 +1,251 @@ +use anyhow::{anyhow, Result}; +use cudarc::{ + driver::{safe::LaunchConfig, CudaModule, CudaStream, PushKernelArg}, + runtime::sys::cudaDeviceProp, +}; +use std::sync::Arc; +use serde_json::{Map, Value}; +use tig_challenges::vector_search::{Challenge, Solution}; + +pub fn solve_challenge( + challenge: &Challenge, + save_solution: &dyn Fn(&Solution) -> anyhow::Result<()>, + hyperparameters: &Option>, + module: Arc, + stream: Arc, + prop: &cudaDeviceProp, +) -> anyhow::Result<()> { + Err(anyhow!("This algorithm is no longer compatible.")) +} + +// Old code that is no longer compatible +#[cfg(none)] +mod dead_code { + // TIG's UI uses the pattern `tig_challenges::` to automatically detect your algorithm's challenge + use anyhow::Result; + use tig_challenges::vector_search::{Challenge, Solution}; + + #[derive(Debug)] + struct KDTreeNode { + point: Vec, + left: Option>, + right: Option>, + } + + #[derive(Debug)] + struct KDTree { + root: Option>, + k: usize, + } + + impl KDTreeNode { + pub fn nearest_neighbor(&self, target: &[f32], best: &mut Option<(Vec, f32)>, depth: usize, k: usize) -> Option<(Vec, f32)> { + let axis = depth % k; + let dist = euclidean_distance(&self.point, target); + + if best.is_none() || dist < best.clone().unwrap().1 { + *best = Some((self.point.clone(), dist)); + } + + let next_branch = if target[axis] < self.point[axis] { + &self.left + } else { + &self.right + }; + + if let Some(branch) = next_branch { + branch.nearest_neighbor(target, best, depth + 1, k); + } + + let other_branch = if target[axis] < self.point[axis] { + &self.right + } else { + &self.left + }; + + if let Some(branch) = other_branch { + if (target[axis] - self.point[axis]).abs() < best.clone().unwrap().1 { + branch.nearest_neighbor(target, best, depth + 1, k); + } + } + + best.clone() + } + } + + impl KDTree { + pub fn new(points: Vec>, depth: usize, k: usize) -> Option> { + if points.is_empty() { + return None; + } + + let axis = depth % k; + let mut sorted_points = points.clone(); + sorted_points.sort_by(|a, b| a[axis].partial_cmp(&b[axis]).unwrap()); + + let median = sorted_points.len() / 2; + + Some(Box::new(KDTreeNode { + point: sorted_points[median].clone(), + left: KDTree::new(sorted_points[..median].to_vec(), depth + 1, k), + right: KDTree::new(sorted_points[median + 1..].to_vec(), depth + 1, k), + })) + } + + pub fn nearest_neighbor(&self, target: &[f32]) -> Option<(Vec, f32)> { + if let Some(root) = &self.root { + let mut best = None; + root.nearest_neighbor(target, &mut best, 0, self.k); + best + } else { + None + } + } + } + + pub fn solve_challenge(challenge: &Challenge) -> Result> { + let k = challenge.vector_database[0].len(); + let tree = KDTree { + root: KDTree::new(challenge.vector_database.clone(), 0, k), + k, + }; + + let mut indexes = Vec::new(); + + for query in &challenge.query_vectors { + if let Some((best_point, _)) = tree.nearest_neighbor(query) { + indexes.push( + challenge.vector_database.iter().position(|x| x == &best_point).unwrap() + ); + } + } + + let solution = Solution { indexes }; + Ok(Some(solution)) + } + + pub fn euclidean_distance(a: &[f32], b: &[f32]) -> f32 { + a.iter() + .zip(b) + .map(|(&x1, &x2)| (x1 - x2) * (x1 - x2)) + .sum::() + .sqrt() + } + + #[cfg(feature = "cuda")] + mod gpu_optimisation { + use super::*; + use cudarc::driver::*; + use std::{collections::HashMap, sync::Arc}; + use tig_challenges::CudaKernel; + + pub const KERNEL: Option = Some(CudaKernel { + src: r#" + extern "C" __global__ void kd_tree_nearest_neighbor( + const float* query_vectors, + const float* vector_database, + int num_vectors, + int vector_size, + int num_queries, + int* closest_indices, + float* min_distances + ) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + int qid = blockIdx.y; + + if (idx < num_vectors && qid < num_queries) { + float distance = 0.0f; + for (int i = 0; i < vector_size; ++i) { + float diff = query_vectors[qid * vector_size + i] - vector_database[idx * vector_size + i]; + distance += diff * diff; + } + distance = sqrtf(distance); + + if (distance < min_distances[qid]) { + min_distances[qid] = distance; + closest_indices[qid] = idx; + } + } + } + "#, + funcs: &["kd_tree_nearest_neighbor"], + }); + + pub fn cuda_solve_challenge( + challenge: &Challenge, + dev: &Arc, + mut funcs: HashMap<&'static str, CudaFunction>, + ) -> anyhow::Result> { + let vector_database: &Vec> = &challenge.vector_database; + let query_vectors: &Vec> = &challenge.query_vectors; + + let num_vectors = vector_database.len(); + let vector_size = vector_database[0].len(); + let num_queries = query_vectors.len(); + + let vector_db_flat: Vec = vector_database.iter().flatten().cloned().collect(); + let query_vectors_flat: Vec = query_vectors.iter().flatten().cloned().collect(); + + let vector_db_dev = dev.htod_sync_copy(&vector_db_flat)?; + let query_dev = dev.htod_sync_copy(&query_vectors_flat)?; + + let closest_indices_dev: CudaSlice = unsafe { dev.alloc(num_queries)? }; + + // Initialize min_distances_dev to std::f32::MAX + let mut min_distances_host = vec![std::f32::MAX; num_queries]; + let min_distances_dev = dev.htod_sync_copy(&min_distances_host)?; + + let stream = dev.fork_default_stream()?; + + let block_dim = (256, 1, 1); + let grid_dim = ( + ((num_vectors + block_dim.0 as usize - 1) / block_dim.0 as usize) as u32, + num_queries as u32, + 1, + ); + + let func = funcs.get_mut("kd_tree_nearest_neighbor").unwrap().clone(); + + unsafe { + func.launch_on_stream( + &stream, + LaunchConfig { + block_dim, + grid_dim, + shared_mem_bytes: 0, + }, + ( + &query_dev, + &vector_db_dev, + num_vectors as i32, + vector_size as i32, + num_queries as i32, + &closest_indices_dev, + &min_distances_dev, + ), + )?; + } + + dev.wait_for(&stream)?; + + let closest_indices = dev.dtoh_sync_copy(&closest_indices_dev)?; + + let indexes: Vec = closest_indices.iter().map(|&i| i as usize).collect(); + + Ok(Some(Solution { indexes })) + } + } + + #[cfg(feature = "cuda")] + pub use gpu_optimisation::{cuda_solve_challenge, KERNEL}; + + + + + + +} + +pub fn help() { + println!("No help information available."); +}