diff --git a/Cargo.toml b/Cargo.toml index b832d19..0212c77 100755 --- a/Cargo.toml +++ b/Cargo.toml @@ -15,4 +15,4 @@ opt-level = "z" codegen-units = 1 [profile.release] -lto = true +debug = true \ No newline at end of file diff --git a/schemsearch-cli/src/main.rs b/schemsearch-cli/src/main.rs index 6a08320..a852c8a 100755 --- a/schemsearch-cli/src/main.rs +++ b/schemsearch-cli/src/main.rs @@ -15,35 +15,35 @@ * along with this program. If not, see . */ -mod types; mod json_output; mod sinks; mod stderr; +mod types; -use std::fmt::Debug; -use std::io::Write; -use clap::{command, Arg, ArgAction, ValueHint}; -use std::path::PathBuf; -use std::str::FromStr; -use clap::error::ErrorKind; +use crate::sinks::{OutputFormat, OutputSink}; +use crate::stderr::MaschineStdErr; +#[cfg(feature = "sql")] +use crate::types::SqlSchematicSupplier; use crate::types::{PathSchematicSupplier, SchematicSupplier, SchematicSupplierType}; +use clap::error::ErrorKind; +use clap::{command, Arg, ArgAction, ValueHint}; #[cfg(feature = "sql")] use futures::executor::block_on; +use indicatif::*; use rayon::prelude::*; use rayon::ThreadPoolBuilder; +use schemsearch_common::{Match, SearchBehavior}; +use schemsearch_files::SpongeSchematic; +use schemsearch_lib::nbt_search::has_invalid_nbt; +use schemsearch_lib::search::search; #[cfg(feature = "sql")] use schemsearch_sql::filter::SchematicFilter; #[cfg(feature = "sql")] use schemsearch_sql::load_all_schematics; -#[cfg(feature = "sql")] -use crate::types::SqlSchematicSupplier; -use indicatif::*; -use schemsearch_common::{Match, SearchBehavior}; -use schemsearch_files::SpongeSchematic; -use crate::sinks::{OutputFormat, OutputSink}; -use crate::stderr::MaschineStdErr; -use schemsearch_lib::nbt_search::has_invalid_nbt; -use schemsearch_lib::search::search; +use std::fmt::Debug; +use std::io::Write; +use std::path::PathBuf; +use std::str::FromStr; fn main() { #[allow(unused_mut)] @@ -179,7 +179,7 @@ fn main() { .bin_name("schemsearch"); #[cfg(feature = "sql")] - let mut cmd = cmd + let mut cmd = cmd .arg( Arg::new("sql") .help("Use the SteamWar SQL Database") @@ -219,7 +219,9 @@ fn main() { ignore_air: matches.get_flag("ignore-air"), air_as_any: matches.get_flag("air-as-any"), ignore_entities: matches.get_flag("ignore-entities"), - threshold: *matches.get_one::("threshold").expect("Couldn't get threshold"), + threshold: *matches + .get_one::("threshold") + .expect("Couldn't get threshold"), invalid_nbt: matches.get_flag("invalid-nbt"), opencl: matches.get_flag("opencl"), }; @@ -228,7 +230,11 @@ fn main() { Some(p) => match SpongeSchematic::load(&PathBuf::from(p)) { Ok(x) => Some(x), Err(e) => { - cmd.error(ErrorKind::Io, format!("Error while loading Pattern: {}", e.to_string())).exit(); + cmd.error( + ErrorKind::Io, + format!("Error while loading Pattern: {}", e.to_string()), + ) + .exit(); } }, None => None, @@ -269,60 +275,103 @@ fn main() { } for schem in block_on(load_all_schematics(filter)) { schematics.push(SchematicSupplierType::SQL(SqlSchematicSupplier { - node: schem + node: schem, })) - }; + } } if schematics.is_empty() { - cmd.error(ErrorKind::MissingRequiredArgument, "No schematics specified").exit(); + cmd.error( + ErrorKind::MissingRequiredArgument, + "No schematics specified", + ) + .exit(); } - let output: Vec<&(OutputFormat, OutputSink)> = matches.get_many::<(OutputFormat, OutputSink)>("output").expect("Error").collect(); - let mut output: Vec<(OutputFormat, Box)> = output.into_iter().map(|x| (x.0.clone(), x.1.output())).collect(); + let output: Vec<&(OutputFormat, OutputSink)> = matches + .get_many::<(OutputFormat, OutputSink)>("output") + .expect("Error") + .collect(); + let mut output: Vec<(OutputFormat, Box)> = output + .into_iter() + .map(|x| (x.0.clone(), x.1.output())) + .collect(); for x in &mut output { - write!(x.1, "{}", x.0.start(schematics.len() as u32, &search_behavior, start.elapsed().as_millis())).unwrap(); + write!( + x.1, + "{}", + x.0.start( + schematics.len() as u32, + &search_behavior, + start.elapsed().as_millis() + ) + ) + .unwrap(); } - ThreadPoolBuilder::new().num_threads(*matches.get_one::("threads").expect("Could not get threads")).build_global().unwrap(); + ThreadPoolBuilder::new() + .num_threads( + *matches + .get_one::("threads") + .expect("Could not get threads"), + ) + .build_global() + .unwrap(); let bar = ProgressBar::new(schematics.len() as u64); // "maschine" - bar.set_style(ProgressStyle::with_template("[{elapsed}, ETA: {eta}] {wide_bar} {pos}/{len} {per_sec}").unwrap()); - let term_size = *matches.get_one::("machine").expect("Could not get machine"); + bar.set_style( + ProgressStyle::with_template("[{elapsed}, ETA: {eta}] {wide_bar} {pos}/{len} {per_sec}") + .unwrap(), + ); + let term_size = *matches + .get_one::("machine") + .expect("Could not get machine"); if term_size != 0 { - bar.set_draw_target(ProgressDrawTarget::term_like(Box::new(MaschineStdErr { size: term_size }))) + bar.set_draw_target(ProgressDrawTarget::term_like(Box::new(MaschineStdErr { + size: term_size, + }))) } - let max_matching = *matches.get_one::("limit").expect("Could not get max-matching"); + let max_matching = *matches + .get_one::("limit") + .expect("Could not get max-matching"); - let matches: Vec = schematics.par_iter().progress_with(bar).map(|schem| { - match schem { + let matches: Vec = schematics + .par_iter() + .progress_with(bar) + .map(|schem| match schem { SchematicSupplierType::PATH(schem) => { let schematic = match load_schem(&schem.path) { Some(x) => x, - None => return SearchResult { - name: schem.get_name(), - matches: Vec::default(), + None => { + return SearchResult { + name: schem.get_name(), + matches: Vec::default(), + } } }; search_in_schem(schematic, pattern.as_ref(), search_behavior, schem) } #[cfg(feature = "sql")] - SchematicSupplierType::SQL(schem) => { - match schem.get_schematic() { - Ok(schematic) => search_in_schem(schematic, pattern.as_ref(), search_behavior, schem), - Err(e) => { - eprintln!("Error while loading schematic ({}): {}", schem.get_name(), e.to_string()); - SearchResult { - name: schem.get_name(), - matches: Vec::default(), - } + SchematicSupplierType::SQL(schem) => match schem.get_schematic() { + Ok(schematic) => { + search_in_schem(schematic, pattern.as_ref(), search_behavior, schem) + } + Err(e) => { + eprintln!( + "Error while loading schematic ({}): {}", + schem.get_name(), + e.to_string() + ); + SearchResult { + name: schem.get_name(), + matches: Vec::default(), } } - } - } - }).collect(); + }, + }) + .collect(); let mut matches_count = 0; @@ -340,14 +389,18 @@ fn main() { } } - let end = std::time::Instant::now(); for x in &mut output { - write!(x.1, "{}", x.0.end(end.duration_since(start).as_millis())).unwrap(); + write!(x.1, "{}", x.0.end(start.elapsed())).unwrap(); x.1.flush().unwrap(); } } -fn search_in_schem(schematic: SpongeSchematic, pattern: Option<&SpongeSchematic>, search_behavior: SearchBehavior, schem: &impl SchematicSupplier) -> SearchResult { +fn search_in_schem( + schematic: SpongeSchematic, + pattern: Option<&SpongeSchematic>, + search_behavior: SearchBehavior, + schem: &impl SchematicSupplier, +) -> SearchResult { if search_behavior.invalid_nbt { if has_invalid_nbt(schematic) { SearchResult { @@ -377,7 +430,11 @@ fn load_schem(schem_path: &PathBuf) -> Option { match SpongeSchematic::load(schem_path) { Ok(x) => Some(x), Err(e) => { - println!("Error while loading schematic ({}): {}", schem_path.to_str().unwrap(), e.to_string()); + println!( + "Error while loading schematic ({}): {}", + schem_path.to_str().unwrap(), + e.to_string() + ); None } } @@ -388,4 +445,3 @@ struct SearchResult { name: String, matches: Vec, } - diff --git a/schemsearch-cli/src/sinks.rs b/schemsearch-cli/src/sinks.rs index e6b52ac..ae99e62 100755 --- a/schemsearch-cli/src/sinks.rs +++ b/schemsearch-cli/src/sinks.rs @@ -1,11 +1,11 @@ -use std::fs::File; -use std::io::BufWriter; -use std::str::FromStr; -use std::io::Write; -use std::time::Duration; +use crate::json_output::{EndEvent, FoundEvent, InitEvent, JsonEvent}; use indicatif::HumanDuration; use schemsearch_common::{Match, SearchBehavior}; -use crate::json_output::{EndEvent, FoundEvent, InitEvent, JsonEvent}; +use std::fs::File; +use std::io::BufWriter; +use std::io::Write; +use std::str::FromStr; +use std::time::Duration; #[derive(Debug, Clone)] pub enum OutputSink { @@ -18,7 +18,7 @@ pub enum OutputSink { pub enum OutputFormat { Text, CSV, - JSON + JSON, } impl FromStr for OutputFormat { @@ -29,7 +29,7 @@ impl FromStr for OutputFormat { "text" => Ok(OutputFormat::Text), "csv" => Ok(OutputFormat::CSV), "json" => Ok(OutputFormat::JSON), - _ => Err(format!("'{}' is not a valid output format", s)) + _ => Err(format!("'{}' is not a valid output format", s)), } } } @@ -41,7 +41,7 @@ impl FromStr for OutputSink { match s { "std" => Ok(OutputSink::Stdout), "err" => Ok(OutputSink::Stderr), - _ => Ok(OutputSink::File(s.to_string())) + _ => Ok(OutputSink::File(s.to_string())), } } } @@ -51,7 +51,7 @@ impl OutputSink { match self { OutputSink::Stdout => Box::new(std::io::stdout()), OutputSink::Stderr => Box::new(std::io::stderr()), - OutputSink::File(path) => Box::new(BufWriter::new(File::create(path).unwrap())) + OutputSink::File(path) => Box::new(BufWriter::new(File::create(path).unwrap())), } } } @@ -59,12 +59,21 @@ impl OutputSink { impl OutputFormat { pub fn found_match(&self, name: &String, pos: Match) -> String { match self { - OutputFormat::Text => format!("Found match in '{}' at x: {}, y: {}, z: {}, % = {}\n", name, pos.x, pos.y, pos.z, pos.percent), - OutputFormat::CSV => format!("{},{},{},{},{}\n", name, pos.x, pos.y, pos.z, pos.percent), - OutputFormat::JSON => format!("{}\n", serde_json::to_string(&JsonEvent::Found(FoundEvent { - name: name.clone(), - match_: pos, - })).unwrap()) + OutputFormat::Text => format!( + "Found match in '{}' at x: {}, y: {}, z: {}, % = {}\n", + name, pos.x, pos.y, pos.z, pos.percent + ), + OutputFormat::CSV => { + format!("{},{},{},{},{}\n", name, pos.x, pos.y, pos.z, pos.percent) + } + OutputFormat::JSON => format!( + "{}\n", + serde_json::to_string(&JsonEvent::Found(FoundEvent { + name: name.clone(), + match_: pos, + })) + .unwrap() + ), } } @@ -72,19 +81,29 @@ impl OutputFormat { match self { OutputFormat::Text => format!("Starting search in {} schematics\n", total), OutputFormat::CSV => "Name,X,Y,Z,Percent\n".to_owned(), - OutputFormat::JSON => format!("{}\n", serde_json::to_string(&JsonEvent::Init(InitEvent { - total, - search_behavior: search_behavior.clone(), - start_time, - })).unwrap()) + OutputFormat::JSON => format!( + "{}\n", + serde_json::to_string(&JsonEvent::Init(InitEvent { + total, + search_behavior: search_behavior.clone(), + start_time, + })) + .unwrap() + ), } } - pub fn end(&self, end_time: u128) -> String { + pub fn end(&self, end_time: Duration) -> String { match self { - OutputFormat::Text => format!("Search complete in {}\n", HumanDuration(Duration::from_millis(end_time as u64))), - OutputFormat::CSV => format!("{}\n", end_time), - OutputFormat::JSON => format!("{}\n", serde_json::to_string(&JsonEvent::End(EndEvent{ end_time })).unwrap()) + OutputFormat::Text => format!("Search complete in {:?}\n", end_time), + OutputFormat::CSV => format!("{:?}\n", end_time), + OutputFormat::JSON => format!( + "{}\n", + serde_json::to_string(&JsonEvent::End(EndEvent { + end_time: end_time.as_millis() + })) + .unwrap() + ), } } -} \ No newline at end of file +} diff --git a/schemsearch-common/src/lib.rs b/schemsearch-common/src/lib.rs index dcae49f..a567cb3 100644 --- a/schemsearch-common/src/lib.rs +++ b/schemsearch-common/src/lib.rs @@ -34,3 +34,14 @@ pub struct Match { pub z: u16, pub percent: f32, } + +#[macro_export] +macro_rules! time { + ($name:ident, $body:block) => {{ + let start = std::time::Instant::now(); + let result = $body; + let duration = start.elapsed(); + println!("{} took {:?}", stringify!($name), duration); + result + }}; +} diff --git a/schemsearch-lib/src/.idea/modules.xml b/schemsearch-lib/src/.idea/modules.xml new file mode 100644 index 0000000..f669a0e --- /dev/null +++ b/schemsearch-lib/src/.idea/modules.xml @@ -0,0 +1,8 @@ + + + + + + + + \ No newline at end of file diff --git a/schemsearch-lib/src/.idea/src.iml b/schemsearch-lib/src/.idea/src.iml new file mode 100644 index 0000000..bc2cd87 --- /dev/null +++ b/schemsearch-lib/src/.idea/src.iml @@ -0,0 +1,8 @@ + + + + + + + + \ No newline at end of file diff --git a/schemsearch-lib/src/.idea/vcs.xml b/schemsearch-lib/src/.idea/vcs.xml new file mode 100644 index 0000000..b2bdec2 --- /dev/null +++ b/schemsearch-lib/src/.idea/vcs.xml @@ -0,0 +1,6 @@ + + + + + + \ No newline at end of file diff --git a/schemsearch-lib/src/.idea/workspace.xml b/schemsearch-lib/src/.idea/workspace.xml new file mode 100644 index 0000000..8234165 --- /dev/null +++ b/schemsearch-lib/src/.idea/workspace.xml @@ -0,0 +1,63 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + 1715303674752 + + + + + + \ No newline at end of file diff --git a/schemsearch-lib/src/search.rs b/schemsearch-lib/src/search.rs index da5f88d..a929d39 100755 --- a/schemsearch-lib/src/search.rs +++ b/schemsearch-lib/src/search.rs @@ -1,16 +1,19 @@ +use crate::pattern_mapper::{match_palette, match_palette_adapt}; use math::round::ceil; -use schemsearch_common::Match; +use schemsearch_common::time; +use schemsearch_common::{Match, SearchBehavior}; use schemsearch_files::SpongeSchematic; use schemsearch_ocl_matcher::ocl_search; -use crate::{SearchBehavior}; -use crate::pattern_mapper::{match_palette, match_palette_adapt}; pub fn search( schem: SpongeSchematic, pattern_schem: &SpongeSchematic, search_behavior: SearchBehavior, ) -> Vec { - if schem.width < pattern_schem.width || schem.height < pattern_schem.height || schem.length < pattern_schem.length { + if schem.width < pattern_schem.width + || schem.height < pattern_schem.height + || schem.length < pattern_schem.length + { return Vec::new(); } @@ -18,17 +21,27 @@ pub fn search( return Vec::new(); } - let pattern_schem = match_palette(&schem, &pattern_schem, search_behavior.ignore_block_data); + let pattern_schem = time!(match_palette, { + match_palette(&schem, &pattern_schem, search_behavior.ignore_block_data) + }); let mut matches: Vec = Vec::with_capacity(4); let schem_data = if search_behavior.ignore_block_data { - match_palette_adapt(&schem, &pattern_schem.palette, search_behavior.ignore_block_data) + match_palette_adapt( + &schem, + &pattern_schem.palette, + search_behavior.ignore_block_data, + ) } else { schem.block_data }; - let air_id = if search_behavior.ignore_air || search_behavior.air_as_any { pattern_schem.palette.get("minecraft:air").unwrap_or(&-1) } else { &-1}; + let air_id = if search_behavior.ignore_air || search_behavior.air_as_any { + pattern_schem.palette.get("minecraft:air").unwrap_or(&-1) + } else { + &-1 + }; let pattern_blocks = pattern_schem.block_data.len() as f32; let i_pattern_blocks = pattern_blocks as i32; @@ -42,29 +55,42 @@ pub fn search( let schem_length = schem.length as usize; if search_behavior.opencl { - return ocl_search(schem_data.as_slice(), [schem_width, schem_height, schem_length], pattern_schem.block_data.as_slice(), [pattern_width, pattern_height, pattern_length], *air_id, search_behavior).unwrap() + return time!(ocl_search, { + ocl_search( + schem_data.as_slice(), + [schem_width, schem_height, schem_length], + pattern_schem.block_data.as_slice(), + [pattern_width, pattern_height, pattern_length], + *air_id, + search_behavior, + ) + .unwrap() + }); } let schem_data = schem_data.as_ptr(); let pattern_data = pattern_schem.block_data.as_ptr(); - let skip_amount = ceil((pattern_blocks * (1.0 - search_behavior.threshold)) as f64, 0) as i32; + let skip_amount = ceil( + (pattern_blocks * (1.0 - search_behavior.threshold)) as f64, + 0, + ) as i32; for y in 0..=schem_height - pattern_height { for z in 0..=schem_length - pattern_length { for x in 0..=schem_width - pattern_width { let mut not_matching = 0; - 'outer: - for j in 0..pattern_height { + 'outer: for j in 0..pattern_height { for k in 0..pattern_length { - 'inner: - for i in 0..pattern_width { + 'inner: for i in 0..pattern_width { let index = (x + i) + schem_width * ((z + k) + (y + j) * schem_length); let pattern_index = i + pattern_width * (k + j * pattern_length); let data = unsafe { *schem_data.add(index) }; let pattern_data = unsafe { *pattern_data.add(pattern_index) }; - if (search_behavior.ignore_air && data != *air_id) || (search_behavior.air_as_any && pattern_data != *air_id) { + if (search_behavior.ignore_air && data != *air_id) + || (search_behavior.air_as_any && pattern_data != *air_id) + { continue 'inner; } if data != pattern_data { @@ -90,4 +116,4 @@ pub fn search( } return matches; -} \ No newline at end of file +} diff --git a/schemsearch-ocl-matcher/src/kernel.cl b/schemsearch-ocl-matcher/src/kernel.cl index 202f22d..55cd8a0 100644 --- a/schemsearch-ocl-matcher/src/kernel.cl +++ b/schemsearch-ocl-matcher/src/kernel.cl @@ -1,51 +1,35 @@ -__kernel void add(__global int* result, - __global uint* schem, - __global uint* pattern, - const int p_width, - const int p_height, - const int p_depth, - const uint air_id, - const int ignore_air, - const int air_as_any, - const int skipamount) { - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); - - int width = get_global_size(0); - int height = get_global_size(1); - int depth = get_global_size(2); - - if (x > width - p_width || y > height - p_height || z > depth - p_depth) { - return; - } - - int wrong_blocks = 0; - for (int py = 0; py < p_height; py++) { - for (int pz = 0; pz < p_depth; pz++) { - for (int px = 0; px < p_width; px++) { - int s_idx = (x + px) + width * ((z + pz) + (y + py) * depth); - int p_idx = px + p_width * (pz + py * p_depth); - - uint schem_block = schem[s_idx]; - uint pattern_block = pattern[p_idx]; +// Use 3d_img +// Weniger Allocs an Buffern +// Pattern Parallelisieren mit Local Workern? +// To Match on GPU +// Weniger Worker, Mehr Parameter! +// Pattern als Kernel Konstante - if ((ignore_air && schem_block != air_id) || (air_as_any && pattern_block != air_id)) { - continue; - } +__kernel void add(__global int *result, __global uint *schem, + __constant uint *pattern, const int width, const int height, + const int depth, const int p_width, const int p_height, + const int p_depth, const uint air_id, const int ignore_air, + const int air_as_any, const int skipamount) { + int x = get_global_id(0); + int y = get_global_id(2); + int z = get_global_id(1); - if (schem_block != pattern_block) { - wrong_blocks++; - if (wrong_blocks > skipamount) { - int idx = x + z * width + y * width * depth; - result[idx] = wrong_blocks; - return; - } - } - } - } + int wrong_blocks = 0; + for (int py = 0; py < p_height; py++) { + for (int pz = 0; pz < p_depth; pz++) { + for (int px = 0; px < p_width; px++) { + // if ((ignore_air && schem_block != air_id) || (air_as_any && + // pattern_block != air_id)) { + // continue; // TODO: PROBLEM! + // } + + wrong_blocks += + schem[(x + px) + width * ((z + pz) + (y + py) * depth)] != + pattern[px + p_width * (pz + py * p_depth)]; + } } - - int idx = x + z * width + y * width * depth; - result[idx] = wrong_blocks; + } + + int idx = x + z * width + y * width * depth; + result[idx] = wrong_blocks; } diff --git a/schemsearch-ocl-matcher/src/lib.rs b/schemsearch-ocl-matcher/src/lib.rs index d52a502..be6e57d 100644 --- a/schemsearch-ocl-matcher/src/lib.rs +++ b/schemsearch-ocl-matcher/src/lib.rs @@ -1,8 +1,9 @@ -use std::sync::OnceLock; -use ocl::{Buffer, MemFlags, ProQue, core}; -use ocl::SpatialDims::Three; -use schemsearch_common::{Match, SearchBehavior}; use math::round::ceil; +use ocl::SpatialDims::Three; +use ocl::{core, Buffer, Image, MemFlags, ProQue}; +use schemsearch_common::{time, Match, SearchBehavior}; +use std::sync::OnceLock; +use std::time; const KERNEL: &str = include_str!("kernel.cl"); @@ -20,7 +21,15 @@ pub fn ocl_search( air_id: i32, search_behavior: SearchBehavior, ) -> Result, String> { - search_ocl(schem, schem_size, pattern, pattern_size, air_id, search_behavior).map_err(|e| e.to_string()) + search_ocl( + schem, + schem_size, + pattern, + pattern_size, + air_id, + search_behavior, + ) + .map_err(|e| e.to_string()) } fn search_ocl( @@ -38,59 +47,79 @@ fn search_ocl( let schem_width = schem_size[0]; let schem_height = schem_size[1]; let schem_length = schem_size[2]; - + let pattern_blocks = (pattern_width * pattern_height * pattern_length) as f32; - - let skip_amount = ceil((pattern_blocks * (1.0 - search_behavior.threshold)) as f64, 0) as i32; + + let skip_amount = ceil( + (pattern_blocks * (1.0 - search_behavior.threshold)) as f64, + 0, + ) as i32; let cell = &PRO_QUEU_CELL; - let mut pro_que = cell.get_or_init(|| { - ProQue::builder() - .src(KERNEL) - .build().unwrap() - }).clone(); + let mut pro_que = time!(get_pro_que, { + cell.get_or_init(|| ProQue::builder().src(KERNEL).build().unwrap()) + .clone() + }); + pro_que.set_dims(Three(schem_width, schem_length, schem_height)); - pro_que.set_dims(Three(schem_width, schem_height, schem_length)); + let buffer = time!(create_result_buffer, { + Buffer::builder() + .queue(pro_que.queue().clone()) + .flags(MemFlags::new().read_write()) + .fill_val(-1) + .len(schem.len()) + .build() + })?; - let buffer = Buffer::builder() - .queue(pro_que.queue().clone()) - .flags(MemFlags::new().read_write()) - .fill_val(-1) - .len(schem.len()) - .build()?; + let schem_buffer = time!(create_schen_buffer, { + create_schem_buffer(schem, &pro_que) + })?; - let schem_buffer = create_schem_buffer(schem, &pro_que)?; + let pattern_buffer = time!(create_pattern_buffer, { + create_schem_buffer(pattern, &pro_que) + })?; - let pattern_buffer = create_schem_buffer(pattern, &pro_que)?; + let kernel = time!(create_kernel, { + pro_que + .kernel_builder("add") + .arg(&buffer) + .arg(&schem_buffer) + .arg(&pattern_buffer) + .arg(schem_width as i32) + .arg(schem_height as i32) + .arg(schem_length as i32) + .arg(pattern_width as i32) + .arg(pattern_height as i32) + .arg(pattern_length as i32) + .arg(air_id) + .arg(search_behavior.ignore_air as u32) + .arg(search_behavior.air_as_any as u32) + .arg(skip_amount) + .build() + })?; - let kernel = pro_que.kernel_builder("add") - .arg(&buffer) - .arg(&schem_buffer) - .arg(&pattern_buffer) - .arg(pattern_width as i32) - .arg(pattern_height as i32) - .arg(pattern_length as i32) - .arg(air_id) // air_id - .arg(search_behavior.ignore_air as u32) // ignore_air - .arg(search_behavior.air_as_any as u32) // air_as_any - .arg(skip_amount) - .build()?; + unsafe { + time!(run_kernel, { kernel.enq() })?; + } - unsafe { kernel.enq()?; } + let mut vec = vec![0; buffer.len()]; + time!(read_buffer, { + buffer.read(&mut vec).enq()?; + }); - let mut vec = vec![0i32; buffer.len()]; - buffer.read(&mut vec).enq()?; - - Ok(vec.into_iter().enumerate().filter(|(_, v)| *v < skip_amount && *v != -1).map(|(i, v)| { - Match { + Ok(vec + .into_iter() + .enumerate() + .filter(|(_, v)| *v < skip_amount && *v != -1) + .map(|(i, v)| Match { x: (i % schem_width) as u16, y: ((i / (schem_width * schem_length)) % schem_height) as u16, z: ((i / schem_width) % schem_length) as u16, percent: (pattern_blocks - v as f32) / pattern_blocks, - } - }).collect()) + }) + .collect()) } fn create_schem_buffer(pattern: &[i32], pro_que: &ProQue) -> ocl::Result> { @@ -98,6 +127,7 @@ fn create_schem_buffer(pattern: &[i32], pro_que: &ProQue) -> ocl::Result