diff --git a/.gitignore b/.gitignore index f2c20d3c..2b9560b3 100644 --- a/.gitignore +++ b/.gitignore @@ -14,6 +14,3 @@ examples/**/*.cu # nix-direnv /.direnv/ /.envrc - -# cuda header -src/shared_mem_config.h \ No newline at end of file diff --git a/build.rs b/build.rs index b20a72c0..31b8e748 100644 --- a/build.rs +++ b/build.rs @@ -9,7 +9,6 @@ fn main() { println!("cargo:rerun-if-changed=src/run.cu"); println!("cargo:rerun-if-changed=src/hvm.cu"); println!("cargo:rerun-if-changed=src/get_shared_mem.cu"); - println!("cargo:rerun-if-changed=src/shared_mem_config.h"); println!("cargo:rustc-link-arg=-rdynamic"); match cc::Build::new() @@ -35,28 +34,31 @@ fn main() { } // Compile get_shared_mem.cu - if let Ok(output) = Command::new("nvcc") + let shared_mem_value = Command::new("nvcc") .args(&["src/get_shared_mem.cu", "-o", "get_shared_mem"]) .output() - .and_then(|_| Command::new("./get_shared_mem").output()) { - if output.status.success() { - let shared_mem_str = String::from_utf8_lossy(&output.stdout).trim().to_string(); - std::fs::write("src/shared_mem_config.h", format!("#define HVM_SHARED_MEM {}", shared_mem_str)) - .expect("Failed to write shared_mem_config.h"); - } else { - println!("cargo:warning=\x1b[1m\x1b[31mWARNING: Failed to get shared memory size. Using default value.\x1b[0m"); - } - } else { - println!("cargo:warning=\x1b[1m\x1b[31mWARNING: Failed to compile or run get_shared_mem.cu. Using default shared memory value.\x1b[0m"); - } + .and_then(|_| Command::new("./get_shared_mem").output()) + .ok() + .and_then(|output| if output.status.success() { + Some(String::from_utf8_lossy(&output.stdout).trim().to_string()) + } else { + None + }) + .unwrap_or_else(|| { + println!("cargo:warning=\x1b[1m\x1b[31mWARNING: Failed to get shared memory size. Using default value.\x1b[0m"); + "0x2000".to_string() + }); // Clean up temporary executable let _ = std::fs::remove_file("get_shared_mem"); + println!("cargo:warning=\x1b[1m\x1b[33mShared memory size set to: {}\x1b[0m", shared_mem_value); + cc::Build::new() .cuda(true) .file("src/run.cu") .define("IO", None) + .define("HVM_SHARED_MEM", Some(shared_mem_value.as_str())) .flag("-diag-suppress=177") // variable was declared but never referenced .flag("-diag-suppress=550") // variable was set but never used .flag("-diag-suppress=20039") // a __host__ function redeclared with __device__, hence treated as a __host__ __device__ function diff --git a/src/hvm.cu b/src/hvm.cu index d0c9fe6d..e8a9b8b3 100644 --- a/src/hvm.cu +++ b/src/hvm.cu @@ -126,10 +126,9 @@ struct RBag { Pair lo_buf[RLEN]; }; -#include "shared_mem_config.h" - +// Default value for shared memory (96KB) #ifndef HVM_SHARED_MEM -#define HVM_SHARED_MEM 0x2000 // Default value +#define HVM_SHARED_MEM 0x2000 #endif // Local Net