CUDA shaders WIP

Cargo.lock

@@ -737,6 +737,24 @@ dependencies = [
  "memchr",
 ]
 
+[[package]]
+name = "cuda-config"
+version = "0.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4ee74643f7430213a1a78320f88649de309b20b80818325575e393f848f79f5d"
+dependencies = [
+ "glob",
+]
+
+[[package]]
+name = "cuda-driver-sys"
+version = "0.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1d4c552cc0de854877d80bcd1f11db75d42be32962d72a6799b88dcca88fffbd"
+dependencies = [
+ "cuda-config",
+]
+
 [[package]]
 name = "d3d12"
 version = "0.6.0"
@@ -873,6 +891,15 @@ dependencies = [
  "instant",
 ]
 
+[[package]]
+name = "find_cuda_helper"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f9f9e65c593dd01ac77daad909ea4ad17f0d6d1776193fc8ea766356177abdad"
+dependencies = [
+ "glob",
+]
+
 [[package]]
 name = "fixedbitset"
 version = "0.4.2"
@@ -2603,6 +2630,8 @@ version = "0.1.0"
 dependencies = [
  "bytemuck",
  "criterion 0.4.0",
+ "cuda-driver-sys",
+ "find_cuda_helper",
  "futures 0.3.28",
  "lazy_static",
  "metal",

sunscreen_math/Cargo.toml

@@ -16,18 +16,21 @@ ocl = { version = "0.19.4", optional = true }
 futures = { version = "0.3.26", optional = true }
 tokio = { version = "1.25.0", features = ["rt"], optional = true }
 wgpu = { version = "0.15.1", optional = true }
+cuda-driver-sys = { version = "0.3.0", optional = true }
 
 [build-dependencies]
 naga = { version = "0.11.0", optional = true, features = ["wgsl-in"]}
 wgpu-core = { version = "0.15.1", optional = true, features = ["vulkan", "wgsl"] }
 ocl = { version = "0.19.4", optional = true }
+find_cuda_helper = { version = "0.2.0", optional = true }
 
 [dev-dependencies]
 criterion = "0.4.0"
 
 [features]
-default = []
+default = ["cuda"]
 nightly-features = []
+cuda = ["dep:find_cuda_helper", "dep:cuda-driver-sys", "gpu"]
 metal = ["dep:metal", "gpu"]
 webgpu = ["dep:wgpu", "dep:tokio", "dep:futures", "dep:naga", "dep:wgpu-core", "dep:bytemuck", "gpu"]
 opencl = ["dep:ocl", "gpu"]

sunscreen_math/build.rs

@@ -1,3 +1,54 @@
+#[cfg(feature = "cuda")]
+fn compile_cuda_shaders() {
+    use std::{path::PathBuf, process::Command};
+
+    use find_cuda_helper::find_cuda_root;
+
+    let outdir = PathBuf::from(std::env::var("OUT_DIR").unwrap());
+    let cuda_root = find_cuda_root().unwrap();
+    let shaders_dir = PathBuf::from(".")
+        .join("src")
+        .join("cuda_impl")
+        .join("shaders");
+
+    let nvcc = cuda_root.join("bin").join("nvcc");
+    let is_cu_file = |file: &std::fs::DirEntry| {
+        file.file_name().to_string_lossy().ends_with(".cu")
+            && file.file_type().unwrap().is_file()
+    };
+    
+    for config in ["test", "release"] {
+        let shaders = std::fs::read_dir(&shaders_dir)
+            .unwrap()
+            .filter_map(Result::ok)
+            .filter(is_cu_file);
+
+        for s in shaders {
+            let filename = s.file_name().to_string_lossy().into_owned();
+            let srcfile = shaders_dir.join(&filename);
+            let outfile = outdir.join(format!("{filename}.o"));
+
+            let c = Command::new(&nvcc)
+                .arg("-c")
+                .arg("-D").arg("CUDA_C")
+                .arg("-o").arg(outfile)
+                .arg(srcfile)
+                .output()
+                .unwrap();
+
+            if !c.status.success() {
+                println!("===STDOUT===");
+                println!("{}", String::from_utf8_lossy(&c.stdout));
+                println!("===STDERR===");
+                println!("{}", String::from_utf8_lossy(&c.stderr));
+                panic!("nvcc compilation failed");
+            }
+        }
+    }
+
+
+}
+
 #[cfg(feature = "opencl")]
 fn compile_opencl_shaders() {
     use std::ffi::CString;
@@ -214,4 +265,7 @@ fn main() {
 
     #[cfg(feature = "opencl")]
     compile_opencl_shaders();
+
+    #[cfg(feature = "cuda")]
+    compile_cuda_shaders();
 }

sunscreen_math/src/opencl_impl/shaders/sunscreen_math.cl

@@ -1145,7 +1145,7 @@ kernel void scalar_sub(
     global u32* c,
     u32 len
 ) {
-    u32 tid = get_local_id(0);
+    u32 tid = get_global_id(0);
 
     if (tid < len) {
         Scalar29 t_a = Scalar29_unpack(a, tid, len);