From 80553cae3c56523c6a0a2355fc52011787486d0b Mon Sep 17 00:00:00 2001
From: Nicolas Patry <patry.nicolas@protonmail.com>
Date: Tue, 15 Apr 2025 19:11:05 +0200
Subject: [PATCH 1/2] Adding a small example crate for end to end testing.

---
 .gitignore                    |  1 +
 examples/simple/Cargo.toml    | 13 ++++++++++++
 examples/simple/build.rs      |  9 +++++++++
 examples/simple/src/cuda.cu   |  6 ++++++
 examples/simple/src/kernel.rs |  1 +
 examples/simple/src/lib.rs    | 29 ++++++++++++++++++++++++++
 flake.lock                    | 27 +++++++++++++++++++++++++
 flake.nix                     | 38 +++++++++++++++++++++++++++++++++++
 src/lib.rs                    |  7 -------
 9 files changed, 124 insertions(+), 7 deletions(-)
 create mode 100644 examples/simple/Cargo.toml
 create mode 100644 examples/simple/build.rs
 create mode 100644 examples/simple/src/cuda.cu
 create mode 100644 examples/simple/src/kernel.rs
 create mode 100644 examples/simple/src/lib.rs
 create mode 100644 flake.lock
 create mode 100644 flake.nix

diff --git a/.gitignore b/.gitignore
index 96ef6c0..3b1b994 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,2 +1,3 @@
 /target
 Cargo.lock
+/examples/**/target/
diff --git a/examples/simple/Cargo.toml b/examples/simple/Cargo.toml
new file mode 100644
index 0000000..3543328
--- /dev/null
+++ b/examples/simple/Cargo.toml
@@ -0,0 +1,13 @@
+[package]
+name = "simple"
+version = "0.1.0"
+edition = "2024"
+
+[lib]
+crate-type = ["dylib"]
+
+[build-dependencies]
+bindgen_cuda = { version = "0.1.5", path = "../.." }
+
+[dev-dependencies]
+cudarc = { version = "0.15.2", default-features = false, features = ["cuda-version-from-build-system", "dynamic-linking", "driver", "runtime"] }
diff --git a/examples/simple/build.rs b/examples/simple/build.rs
new file mode 100644
index 0000000..5ad6416
--- /dev/null
+++ b/examples/simple/build.rs
@@ -0,0 +1,9 @@
+fn main() {
+    let builder = bindgen_cuda::Builder::default();
+    let bindings = builder.build_ptx().unwrap();
+    bindings.write("src/kernel.rs").unwrap();
+    println!(
+        "cargo:rustc-link-search=native={}",
+        "/run/opengl-driver/lib"
+    );
+}
diff --git a/examples/simple/src/cuda.cu b/examples/simple/src/cuda.cu
new file mode 100644
index 0000000..eb51677
--- /dev/null
+++ b/examples/simple/src/cuda.cu
@@ -0,0 +1,6 @@
+extern "C" __global__ void sin_kernel(float *out, const float *inp, const size_t numel) {
+    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < numel) {
+        out[i] = sin(inp[i]);
+    }
+}
diff --git a/examples/simple/src/kernel.rs b/examples/simple/src/kernel.rs
new file mode 100644
index 0000000..4784915
--- /dev/null
+++ b/examples/simple/src/kernel.rs
@@ -0,0 +1 @@
+pub const CUDA: &str = include_str!(concat!(env!("OUT_DIR"), "/cuda.ptx"));
diff --git a/examples/simple/src/lib.rs b/examples/simple/src/lib.rs
new file mode 100644
index 0000000..a471fb1
--- /dev/null
+++ b/examples/simple/src/lib.rs
@@ -0,0 +1,29 @@
+mod kernel;
+
+#[cfg(test)]
+mod tests {
+    use cudarc::driver::{DriverError, LaunchConfig, PushKernelArg};
+
+    #[test]
+    fn test_simple() -> Result<(), DriverError> {
+        // Get a stream for GPU 0
+        let ctx = cudarc::driver::CudaContext::new(0)?;
+        let stream = ctx.default_stream();
+
+        // copy a rust slice to the device
+        let inp = stream.memcpy_stod(&[1.0f32; 100])?;
+
+        // or allocate directly
+        let mut out = stream.alloc_zeros::<f32>(100)?;
+        // Dynamically load it into the device
+        let ptx = cudarc::nvrtc::Ptx::from_src(crate::kernel::CUDA);
+        let module = ctx.load_module(ptx)?;
+        let sin_kernel = module.load_function("sin_kernel")?;
+        let mut builder = stream.launch_builder(&sin_kernel);
+        builder.arg(&mut out);
+        builder.arg(&inp);
+        builder.arg(&100usize);
+        unsafe { builder.launch(LaunchConfig::for_num_elems(100)) }?;
+        Ok(())
+    }
+}
diff --git a/flake.lock b/flake.lock
new file mode 100644
index 0000000..a44ef7a
--- /dev/null
+++ b/flake.lock
@@ -0,0 +1,27 @@
+{
+  "nodes": {
+    "nixpkgs": {
+      "locked": {
+        "lastModified": 1730531603,
+        "narHash": "sha256-Dqg6si5CqIzm87sp57j5nTaeBbWhHFaVyG7V6L8k3lY=",
+        "owner": "NixOS",
+        "repo": "nixpkgs",
+        "rev": "7ffd9ae656aec493492b44d0ddfb28e79a1ea25d",
+        "type": "github"
+      },
+      "original": {
+        "owner": "NixOS",
+        "ref": "nixos-unstable",
+        "repo": "nixpkgs",
+        "type": "github"
+      }
+    },
+    "root": {
+      "inputs": {
+        "nixpkgs": "nixpkgs"
+      }
+    }
+  },
+  "root": "root",
+  "version": 7
+}
diff --git a/flake.nix b/flake.nix
new file mode 100644
index 0000000..f0b0d3d
--- /dev/null
+++ b/flake.nix
@@ -0,0 +1,38 @@
+{
+  inputs = {
+    nixpkgs.url = "github:NixOS/nixpkgs/nixos-unstable";
+  };
+
+  outputs =
+    { nixpkgs, ... }:
+    let
+      forAllSystems = nixpkgs.lib.genAttrs [
+        "aarch64-linux"
+        "x86_64-linux"
+        "aarch64-darwin"
+      ];
+    in
+    {
+      devShells = forAllSystems (
+        system:
+        let
+          pkgs = import nixpkgs {
+            inherit system;
+            config.allowUnfree = true;
+          };
+        in
+        with pkgs;
+        {
+          default = pkgs.mkShell {
+            nativeBuildInputs = [ pkg-config ];
+            buildInputs = [
+              rustup
+              cudaPackages.cudatoolkit
+              cudaPackages.cuda_nvcc
+            ];
+          };
+
+        }
+      );
+    };
+}
diff --git a/src/lib.rs b/src/lib.rs
index 6078c5a..0695c19 100644
--- a/src/lib.rs
+++ b/src/lib.rs
@@ -325,7 +325,6 @@ impl Builder {
         include_paths.sort();
         include_paths.dedup();
 
-        #[allow(unused)]
         let mut include_options: Vec<String> = include_paths
             .into_iter()
             .map(|s| {
@@ -441,7 +440,6 @@ fn cuda_include_dir() -> Option<PathBuf> {
         "CUDA_TOOLKIT_ROOT_DIR",
         "CUDNN_LIB",
     ];
-    #[allow(unused)]
     let env_vars = env_vars
         .into_iter()
         .map(std::env::var)
@@ -459,13 +457,8 @@ fn cuda_include_dir() -> Option<PathBuf> {
 
     println!("cargo:info={roots:?}");
 
-    #[allow(unused)]
     let roots = roots.into_iter().map(Into::<PathBuf>::into);
 
-    #[cfg(feature = "ci-check")]
-    let root: PathBuf = "ci".into();
-
-    #[cfg(not(feature = "ci-check"))]
     env_vars
         .chain(roots)
         .find(|path| path.join("include").join("cuda.h").is_file())

From a8eb5be01a3f3604e42d4a1d7878843c8eba9230 Mon Sep 17 00:00:00 2001
From: Nicolas Patry <patry.nicolas@protonmail.com>
Date: Tue, 15 Apr 2025 20:38:10 +0200
Subject: [PATCH 2/2] Adding both examples in directory with tested values.

---
 examples/{simple => simple_ptx}/Cargo.toml    |   0
 examples/{simple => simple_ptx}/build.rs      |   4 --
 examples/{simple => simple_ptx}/src/cuda.cu   |   0
 examples/{simple => simple_ptx}/src/kernel.rs |   0
 examples/{simple => simple_ptx}/src/lib.rs    |  12 +++++-
 examples/simple_static/Cargo.toml             |  13 ++++++
 examples/simple_static/build.rs               |   6 +++
 examples/simple_static/libsin.a               | Bin 0 -> 12514 bytes
 examples/simple_static/src/cuda.cu            |  14 +++++++
 examples/simple_static/src/lib.rs             |  39 ++++++++++++++++++
 flake.nix                                     |   2 +
 11 files changed, 85 insertions(+), 5 deletions(-)
 rename examples/{simple => simple_ptx}/Cargo.toml (100%)
 rename examples/{simple => simple_ptx}/build.rs (61%)
 rename examples/{simple => simple_ptx}/src/cuda.cu (100%)
 rename examples/{simple => simple_ptx}/src/kernel.rs (100%)
 rename examples/{simple => simple_ptx}/src/lib.rs (61%)
 create mode 100644 examples/simple_static/Cargo.toml
 create mode 100644 examples/simple_static/build.rs
 create mode 100644 examples/simple_static/libsin.a
 create mode 100644 examples/simple_static/src/cuda.cu
 create mode 100644 examples/simple_static/src/lib.rs

diff --git a/examples/simple/Cargo.toml b/examples/simple_ptx/Cargo.toml
similarity index 100%
rename from examples/simple/Cargo.toml
rename to examples/simple_ptx/Cargo.toml
diff --git a/examples/simple/build.rs b/examples/simple_ptx/build.rs
similarity index 61%
rename from examples/simple/build.rs
rename to examples/simple_ptx/build.rs
index 5ad6416..53d842a 100644
--- a/examples/simple/build.rs
+++ b/examples/simple_ptx/build.rs
@@ -2,8 +2,4 @@ fn main() {
     let builder = bindgen_cuda::Builder::default();
     let bindings = builder.build_ptx().unwrap();
     bindings.write("src/kernel.rs").unwrap();
-    println!(
-        "cargo:rustc-link-search=native={}",
-        "/run/opengl-driver/lib"
-    );
 }
diff --git a/examples/simple/src/cuda.cu b/examples/simple_ptx/src/cuda.cu
similarity index 100%
rename from examples/simple/src/cuda.cu
rename to examples/simple_ptx/src/cuda.cu
diff --git a/examples/simple/src/kernel.rs b/examples/simple_ptx/src/kernel.rs
similarity index 100%
rename from examples/simple/src/kernel.rs
rename to examples/simple_ptx/src/kernel.rs
diff --git a/examples/simple/src/lib.rs b/examples/simple_ptx/src/lib.rs
similarity index 61%
rename from examples/simple/src/lib.rs
rename to examples/simple_ptx/src/lib.rs
index a471fb1..a1c4108 100644
--- a/examples/simple/src/lib.rs
+++ b/examples/simple_ptx/src/lib.rs
@@ -6,12 +6,13 @@ mod tests {
 
     #[test]
     fn test_simple() -> Result<(), DriverError> {
+        let data: Vec<f32> = (0..100).map(|u| u as f32).collect();
         // Get a stream for GPU 0
         let ctx = cudarc::driver::CudaContext::new(0)?;
         let stream = ctx.default_stream();
 
         // copy a rust slice to the device
-        let inp = stream.memcpy_stod(&[1.0f32; 100])?;
+        let inp = stream.memcpy_stod(&data)?;
 
         // or allocate directly
         let mut out = stream.alloc_zeros::<f32>(100)?;
@@ -24,6 +25,15 @@ mod tests {
         builder.arg(&inp);
         builder.arg(&100usize);
         unsafe { builder.launch(LaunchConfig::for_num_elems(100)) }?;
+        let out_host: Vec<f32> = stream.memcpy_dtov(&out)?;
+
+        assert_eq!(out_host.len(), data.len());
+        // Only approximations can be asserted
+        let expected: Vec<_> = data.into_iter().map(f32::sin).collect();
+        for (i, (l, r)) in out_host.into_iter().zip(expected.into_iter()).enumerate() {
+            let diff = (l - r).abs() / (l + 1e-10);
+            assert!(diff < 1e-3, "{l} != {r} (diff = {diff:?}, location = {i})");
+        }
         Ok(())
     }
 }
diff --git a/examples/simple_static/Cargo.toml b/examples/simple_static/Cargo.toml
new file mode 100644
index 0000000..2d5128f
--- /dev/null
+++ b/examples/simple_static/Cargo.toml
@@ -0,0 +1,13 @@
+[package]
+name = "simple"
+version = "0.1.0"
+edition = "2024"
+
+[lib]
+crate-type = ["dylib"]
+
+[build-dependencies]
+bindgen_cuda = { version = "0.1.5", path = "../.." }
+
+[dependencies]
+cudarc = { version = "0.13", default-features = false, features = ["cuda-version-from-build-system", "dynamic-linking", "driver", "runtime"] }
diff --git a/examples/simple_static/build.rs b/examples/simple_static/build.rs
new file mode 100644
index 0000000..2ed4abf
--- /dev/null
+++ b/examples/simple_static/build.rs
@@ -0,0 +1,6 @@
+fn main() {
+    let builder = bindgen_cuda::Builder::default();
+    builder.build_lib("libsin.a");
+    println!("cargo:rustc-link-lib=sin");
+    println!("cargo:rustc-link-search=native={}", ".");
+}
diff --git a/examples/simple_static/libsin.a b/examples/simple_static/libsin.a
new file mode 100644
index 0000000000000000000000000000000000000000..06e86520960f5d882283dd891ab3fb2a33612e68
GIT binary patch
literal 12514
zcmeG?Yj7LKdEY$<fDlE16eW_9EFB5ZB4r4|9S$#%miQKAQ8H~wi5_MkkN_y!BnTfs
z6;kAwj_fq7tk~nJ$9B@0x|*c%Bxxs3Gp(I;95!*9cj`yldhBHCc$CCxoK$(#u^Zdy
z_wC(+AV_7}^heW~?vS_pJ@?yhzx{4+v28s`W8&0q=T2pL2~>b)yYi_hQJpiKbOL~c
z|3BmZiQ_>%kw~RaPfer~`NG_IB5_;|<fpQU+tNlhJw2Kk-Jh8<pRk06>Ev8?;#2}L
z8tkk#h<A6DS_^Z+N$12|D!C&RNu@)2DjXXRWG14)kbick)48*9SSU~lR*o;Mo5NEz
zO*Ivtbly2MGAs%L4$>?F?1m+XV*#Sum)gCo*$W$BJ;X0uPtcF);o8~@@w;Ao7@P5h
zbKaKt#Sw2He({_)8eb@Q-SLYBZ+zjfw=RD1u(vZ6zaV#XGS<7jAH-T+j9(n|wtQvK
z+e-#y;;t;ixDhkTF<&K~y7<D~-g)GSU#P(+WL(pK=js~t=Hd&3)YqUl{wm?`S)5Z#
z1vJ65(QJ-*j}BfK^q!1gKx-}>_KuERyw^LA+)U)?@%RPPoK*Ze)4sWu_+3xanAAn}
zIrVGm>uPC8T^v%s_9Z5f$@pDc??$pYe&&b@T;Jh+3%^2U6p24UqoDsI@rAdaq{g9z
z=i}#JQ_v|*@r5sv<t_0G=e$j5+h%JBg=^yHpH|ex6Jve1RP|5(UU+e2;fL{se`DNN
z<pdc0>5<O~;D*I5*ew<j1R;*~8|xWZe%R5#V`T`Sp`yK7(~6&<-6D$c6r8JIY_asW
zSc;C-9D`nnf{af$Im4enl`j~D<T$o?eP+glS)ItA%BG6o&z|;AWizwjFQkhFf5igh
zlxkf5RMJSNs<D&Pv*XEWR5ev#Iwxkc`9d;V(3e44wHRj$EF&1{Y3hA-Vk$j8H<`#7
z$(c0xt2nVY!c?;)mt#nEa(a5wNajv4^K!;)VYYB4hYFEK8j=%*soCuRt{W2x%1uTg
zF{RDsrV^QDp8Rcot#Zn8-~L|swIa+af_jyvhByi8S5Rgc-n9aM2Vw+R!nar)A7cD#
zR`CA`!`H3AKgaNESKxob@bU`$4qS2)@YT9R-pBBcC4|*Fiugu%#klCbh^2(ZYc^S%
zU<u@NRJ-c#nxF3kN9nHlP6un6{JO=zqOx?#WUkzV4^x4xi)B`Mrvr$M(uv$?wuvY=
zRJ1Gc5t-#7+35fq@|_f#Kx%CQx!Lrgh}j3OG?ZZhaO_vw$}s6BSlUzqd9wpX5e(E~
ze_NJ(N@FFqcdYttt9}dXJQ2*&&;d1J4$z2bTQO)M2)2gLFs(ZV5C5j^&2EeQ|JCog
z2fU|+5=~$UK4gc>E@(Yt#cT%}-dT?@%{G=c{P)&__pHbmBoX<)z7^`u;;Z-;fZWL|
z61?bw0~ljwnU~;SUEnR$lqlku{#RV^<XMd7VMI{S{Jjf0PRk-32dJU`zwUy&&VWqu
z4y{}2;8HydoxyiMTkfb2<&~lY=ihAK1JDcb%8l}xBEbCKZ|{Y1>>uT8@)C<rP!(qH
zkYP9BCWO&mPtj(7(e<@2U=ILabusu<+oiu^eKOc#|FZtYk24ti+q(d;$N+OZ)F1YE
zCp>44m+12q^wy$OY9^Kqn0Qd$Tg1_0laK%X=VU<RsW=$jSj!e6r#BWQi56vP5&c2(
zhO>(9Aa~A}-s1GeWpp>n&*S<ny<B4QPo8BtQQ-b=E!LK3(<zaCIN#p)*-FHL{ku%q
zUZN>vFq@AW8ynsu=6tw|PA5tN*MEP}!3^5-^3PplZ=C^+WPh$5{d3tR*A=0W49>Ox
z2hE>+zfrq~^zXTRCmTP}ssCL28)bgVaJlx^%P?j-gnuQ$)2JW$DZ}R4f5GVPVk6Vf
z?ddS&8f!clPd9?wkY#$ilwN+@PV?<Px3-4G<K;WoF#n2cnEzyozgJvt)bFJ6=5BtA
z#hdr$x+VH$m)qN4zZ}+F%i(tSa=5W!3C8~SPT}NH!Y!R=>%(1?YpK7|_s#hlIwOF>
zmK?_0t1N!lYlX!J&IkHeuJSLVSNKb=@~3S26a8n-hiv*2{fGLzVhyuzjX8euFMETr
z`NR40Y9_tT@{g$5fW6)6)lc!+Z)5TnoypT)fZ0Rx{ib~Vl7HpRjE?h7^6jQPqf<P#
zoBS@~$NBa)nsf#ez0st*Eq%0JwvzuA-*2uLYdj4^$N3Vq8eqE%Zm($SWAda=GhvIr
zJf5|@f5XAX$MO^AtG$ol*k}QOnGf8>+FBM*G=HeiU96E5vd4m34ZDM;vqY2pip$~U
z{HFcTYDl&+9pzgM8(W_?YklH;*Ouq|X%_#iKOX-$U$y3Zm*=;(V*OzJJ8YCb+4^Af
zEmzNXtr-uvUTMA>O?t(A)|S^B(c4YBHJ>!!wdQ<d|1_VPNoVu*48o!bm;G(&FZ<il
z$NY`<wHh)Tzf9vNdy6)i+|AxM=$}4ZA7(NA-&(ZW=!N94A)f3j+H7n-EdMsLe2w{%
z;)DE)`I+LyPWo{@qJLSw&E3rM6*L%<IbZJk?C?)6XcX-@vFN`bizkF_Y%0qzIiK=z
zk44YRY`+GZ1MK~S`Z1SMK4JdFe2&{qge52PWB=YpiNRpm+j*9j^E`53Y1#sA)JNk1
zH|g5~4#UA2WO!vO?8Wz^LUb!06^vH~viBqEH%$F3e=z@<RL<X6EJ@^1YyHUdzP@~q
zqlfY1d@2~f<`8?$^<J=*9{a@h-JDMH<p0v*yt)1ga_LatfhRfLQLL%?{3o{KA<Mi>
z>)BD1aKV+H0s&!~-zId=r@j$k{OAu)lWhN~{jWdAti=8uMcUZT?_u(|{yc_kCx5Vm
z6U=8GqlSGAutDSaJch*REQO+c!&%&497X%wIR8u@BqmE8FnWzh2L`34``CWO!v;J5
z1*Q+}`vC4oh}G8Wt|R^f=x?)_@kjo@Lg}3H`JxM2%5>!VKLGdLKhX2YgX1qPeoz1R
zXWzZ;t*-Th2eaqjyu?TOyVev_Y!0dv^jh#<VyGxpAbZEGCjW;P<3yD}6$tB4H^pUj
zxndz5mH>gY6*XfVfN~j?7xBMZC*kp}#Y)bs%Fk-vuGQJ#6pJSO7PBPz7%g;J-+(3G
zTgI4;hb_jF7M|qk>$X~cGa7_{)n>vQ7Nb%jPfae*{SZZAe9y0tZ?P7S+mh!Mm*e%1
zkU;yqUd~|EomKK=KkX;D9M}6M;#QN7t|CviSIe8{6DwmNZpm|g6lJ9vzJOSo(|(g&
zevYmpzncF)dKLMjmVAMfTV*wQI*%b1w)f_YRI_g(tgVl3@*g(W&tqL#=-h_niN35@
z=a@XsIgQUhhpAsOm7jldJg;B28ocKIqdrNF``KbES5L6bBLJ@}<=9xSa$w8t<#nap
z?`EeU97U@-Ek)O29Y<KkNIF)w7SE5B3ILnCd4O8z(dA_5aqiqH!Idpe|LL@m$D>Up
z>erqALeiK_7nJ-=A{ul0lc|)E&gbzrm9uFjtUFb>qWR5(UB#cAo|?=my5b*?D%B^G
zCp3Lbxn6m1IOJ1QUH2*B2>u5U2n8GAiX-MzVo{%>sn}3s8n!gmr-W+Ziln267C?Ly
zMS@ss&7kP29~Myq%0%#ARS$qWxB=h*jH+Snj_z;5uQyEPr^ct#ia(t#7-y8F1I13K
zKbOSAVZ}ce)?1(<u>tlZ&=gqHtSaCSY6|?N=InC5*Xf_0o!BS{>)|(R@b{<OY$0$$
zi;i_W@c?)dhZ)wDjzgm^Xc}?DXLESut#ssiRBiV<crb%nIx;=MXd}4C8)4Y!3B|ks
zd*N6n)f3!(41Ng)Hq_nSPN#Eb_OyR)8aDa<6#BcJ)2UZM^uUaf!m*zC8TbQZU9eLP
zYzLuL#m8U3R;>%fISmcL`8Wvce2+qVp!;zU@$jk{rm}^pl)u;vcc2?PCJM<G7z|>|
z54R^%I3j<(6}+l14B{d9qM-%5o%wVj=budXKq#mBI^mZMfo^ASM^077jilloNhJFF
z0|`|}QzuRrlK#SMgH-c*z$olZse$e`JTd_3K4?g_!e`ZP=gi!+{|}~8=@#&)!5L`O
zWno=@A9U)zQD_L>AV||7I>la`-V>XkT^~DwpY&M!Hp8VrCOE8V!?9@a0tAfglpu)P
zKs7?qZP29S$5%6miqMx)L%uc;4+zq`z#YB`p6?$Xhy{iR`?{T(nPdpy8t_JZ5fHi{
z){#+rpds1==YoCo;~Tb`V?7}LF|<T*YZnkd8OZcS2DG3)5V{_&4dSCs3)Dd>hRz+&
zD<Itq!uJGix-i!aT24bN9|&ah$iR>uiCz|3wabErtWof!Z;!zGU<7(2!_j~MLIhe7
z`dGBviC*;2WvAA`5c=A=6<T(E7lZ?_4!OiH3w7Ux26VdP1ra_cYzbu0z6>-5Gazo2
zYib0~DWe}WE#$-eaSVdVY^n=_TG)3DK0~2<MWGcux^D~g?ohSpK4_ZOU|nP%9KwEQ
zjrHIO`8L86J2X{~=+R(UkG>Du!@fhXF7%HeMy2b5fuBj&hC&S#ao1}x->o($&S$R}
zF;c2ic*Cae!o^}FCufpMdv<m%Ys^khx9^uc*@+yqYtb7eO^b~M<l2u)o`}x_(#sG`
z1q0s!+ZRCUgz=z;(N$P$6M8{-7kHv$ZBY9I37!YZ4$gNZ@iq|b(4t=u@YG4Pf&9;6
zdp5lZyg7X%G=v`#TlI%TcwAolARJMz!xZvu5cUZ|n!+uZpBq07o|ta~L_3TyZd+c5
z7OEbwix?-{VcpzMWN{0I<@Y7O5qwP+PC$D~wP7?oD%d^=106;v{1e$98`}&`k<Adl
zZihbhjObCdTg6aN?f!-&{1F6mn0!F|lGGaV37$yzHoGtko}l);VEX~qG%17Ncu5Qf
zb>9O3*MeiaBz#*Ez6zoa_UD9-ux}S;l#nh6FUVVrHnCX^3Olqwv#>506k?h>b_vjN
z*TC9Kup#=lu;vn68$prg=$N!EvMVcig0VW;^=nD?>VcEc84MJ`tE(G@W^G6g>6%Z6
zE?te2DGF>1w%{^}6W3D!H#za_JB5ojo1WYVf=%%0p=(4>*!K?rkHL}jbk0AHpEx(-
z#~Ju?;qy{c<Te<_m0$cO`r{^WZW3%KAf`t#ICpjAA__nc!1H!#6WDHna46vO3oWsu
zf;Xh*a2=@7q2W0EdhBBkaVxl<a)8h$6hasP<C&==&IYX@NFfkBAp0fxh@(BMKQ1<9
zJ|Q;huZYrPPLJl>1g}ehRIq;-{W9Di(sf-MRAYnW@0}q0O6<v`w8sTc)Hf(J2L45m
z-eVWWz@w=*2#$Bz#eE>OK}+sgu_gAbC_X5|>%wOBb_x&$S_8-!@r?t#w^nR-z&LD*
zy=d=<`WDbxq1PPnko0GPOn*!t)bV}vVYxN-VcD&|WV81=gxlbQ{eAu6KxAm}258W3
zfRn+<K!4xBafdTpi|)jrzZbS?Az6s3>aOo#`iR^Bw+{rO(a?}y2l+v*udlDa|0SqB
z0m9$dwq=4RpexdS5E^2i0ryZpwD(2ppkWw)lR8L04cnx{ARdI4*bnf%ZgcFZb@sq|
z4YTz_IG#4>i1;Qz+yUDn-S439^>>7|vtm$HDI@N-ZBnnPYYl!(Xbb8;f*Q@~x(7WT
z@!{*@F|ge$><sRjtdrq40zA4_`h_IjC*c65>NaXV+qMYCQkxcj)M4LWCv@7{B4c-A
zj=U3|!4zPm3z)g?aMIRo={>r91A#xMJK`5Hr&>Sg-wOESd*2iEfq5&$FS1{w;}<YH
z<FFRKf~bXO=|Ni{^9F!*S%7{@f0c2W$`E$^nBc1N_`o>&Y-J%F!|xp!*rjytJ$R%`
zQG@=VKVb6Geo?`H>KKXmttBiHj_bnI&EQ(ZQK2&=D#F>DK`6R~O?7t1{m4w)DLN0N
zV*av)dopZW-N4U@u($qELA+T1n6SV8(?SyKK13XM);n%mL-Mx}4;JPGtZ{*rKU_XQ
zg#LQ>z2f2emJdk#>Xq~IaDDH%vtAjfZ|SRd_t!fPt+}ZJGW-8D$ze&-W3b+Ru)bw~
zy>h6&bErP>?t1rVJ?<BOTOmKl@N?@dW;v%+_EI&TJh-w*<&iDXcCaQN@e)3YO`MA|
zg8P*Av;tbhOQrd%{hms{|B2;~4a*OR!}N@LtW7+&F?E-VRDtKnkF&;F4wmarvDS*O
z8+aV?G^f?&8;EGco;i)h1>)Yoe~zn^S!cyoC63x%h-Z2%7jX~jSK+_mD)<wY9>vnb
z;tcgqUWNbsRq!9V3jQ-!!GG&2cqq)|ikU(p5C~+{R4hadR(6Eh?FK(xamb`q|3oea
znPg#nDm$1gBq4EpL<^gr8Ar0XAD_w>(nkMOHffwWQka@P%@_hE!*H@NFo%M1+&)dG
z83k{Y<qJ~@ETsQtOTKT~6jVcr1okz^?wlRQYcJp*83@orh3Y<0x!Eo1{cl`UGu4dF
z<_7Sh$-rzjGc`G9B-x#k#NLq``};-`Hx3UU8akXf+}A%cl&DO+Wbz2R-L{|IfJ@{G
z`Lk#k-lv(kEirNGHZ*=}df5dytGl`Usa5)0=7WmyEXU(*I@7%6Xg>U*YzmTiMF_9=
zP~qS2bWTNW`fZ-dJhrJ2p4K?WKTJq0w9;u^aJ;z=0k$Kqu>w!)iVE@5nHlGQn~+!t
zPq~}pDekEd9zPVz_=O65rGB1gh@Z|mxSoaze)?&K<N5Ihjf=7j$MZZ!dg%96j#scv
zh4?AUaQrqxVyPZ4<wq*j{$y`l#FKu?X<W}Bwy9L}A2893zZKh@|FDH8e){M5F$+)n
zH&);+G-kG?Gc3Y!zY(5fcu}kd*U#~K6`tqBlht_3?`CwFHj*x$$39!Z&uw5Rf1clJ
z<G0+(_ulXt6n@UK<(;CO;vdiF=?)dGeM}d4Rcw~<bU#k@tL)~O$&VPkb~Tr?P;{p)
zku;3t84Hm;ZLXq2>B$*9ekd!N$QPo3x7KD%Y<9Mg_TwgG2V1z<Gr^PDIsD|B&83Y3
zSC>9z-Db1K_&=!>6ed?zzrlJDN9E5gKbz3Xchn{%(7g6?2CIJgTC;J*f@w{XDn1t!
zn^cI6@H{^dj>-?K`Zcj8FK<)r!++9QE&r+|&-)`@D%G%xIAW#PsVrRn8S9=Vr&Y_-
z9*yKR&S2F~C}v}6hhf5eSiBxZG_|Y!|BPbtC(Ww-<m-VW_G4o;|370Tj!OSiPU7}+
z|4$%}@>#Y2qHT!8!ud^T3Fi3sAYrxs7duS}-Ve<`?V-6G$J6?y@mKfXYxU3hkcAa4
z&*52Ytk(YvtMpHE!R5K$dk{zcQ?9QpNB@iCe*lF@K2Rx)?UR-~jk~fCFBM*Yq@v9T
z)+2ApSK_G6<v9M6*r0i>j^CmsUu_p3KbQX!Hb`f+{9~5<3!KratL5=%VO9TMvHV}O
z<j4;V_|N6J|5xjqMCRM+3GhM7uqJF%oL0+IOjXAYqQKEXPk`@G2Nk6XWHzlPf2zwA
M`c{Q3HLK<S6a6ReK>z>%

literal 0
HcmV?d00001

diff --git a/examples/simple_static/src/cuda.cu b/examples/simple_static/src/cuda.cu
new file mode 100644
index 0000000..4673af5
--- /dev/null
+++ b/examples/simple_static/src/cuda.cu
@@ -0,0 +1,14 @@
+#include "cuda.h"
+#include "stdio.h"
+
+extern "C" __global__ void sin_kernel(float *out, const float *inp, const int32_t numel) {
+    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < numel) {
+        out[i] = sin(inp[i]);
+    }
+}
+
+extern "C" int launch_sin2(float *out, const float *inp, int32_t n,  cudaStream_t stream) {
+    sin_kernel<<<n, 1, 0>>>(out, inp, n); 
+    return 0;
+}
diff --git a/examples/simple_static/src/lib.rs b/examples/simple_static/src/lib.rs
new file mode 100644
index 0000000..63a7255
--- /dev/null
+++ b/examples/simple_static/src/lib.rs
@@ -0,0 +1,39 @@
+use std::ffi::{c_int, c_void};
+
+use cudarc::driver::sys::CUstream_st;
+
+unsafe extern "C" {
+    pub fn launch_sin2(out: *mut c_void, inp: *const c_void, n: c_int, stream: &*mut CUstream_st);
+}
+
+#[cfg(test)]
+mod tests {
+    use cudarc::driver::{DevicePtr, DevicePtrMut, DriverError};
+
+    #[test]
+    fn test_simple() -> Result<(), DriverError> {
+        let data: Vec<f32> = (0..100).map(|u| u as f32).collect();
+        // Get a stream for GPU 0
+        let dev = cudarc::driver::CudaDevice::new(0)?;
+
+        // copy a rust slice to the device
+        let inp = dev.htod_copy(data.clone())?;
+        let mut out = dev.alloc_zeros::<f32>(100)?;
+
+        let out_ptr = *out.device_ptr_mut() as *mut core::ffi::c_void;
+        let inp_ptr = *inp.device_ptr() as *const core::ffi::c_void;
+        unsafe { super::launch_sin2(out_ptr, inp_ptr, 100, dev.cu_stream()) };
+
+        let out_host: Vec<f32> = dev.dtoh_sync_copy(&out)?;
+        assert_eq!(out_host.len(), data.len());
+        // Only approximations can be asserted
+        let expected: Vec<_> = data.into_iter().map(f32::sin).collect();
+        println!("Expect {expected:?}");
+        println!("Got {out_host:?}");
+        for (i, (l, r)) in out_host.into_iter().zip(expected.into_iter()).enumerate() {
+            let diff = (l - r).abs() / (l + 1e-10);
+            assert!(diff < 1e-3, "{l} != {r} (diff = {diff:?}, location = {i})");
+        }
+        Ok(())
+    }
+}
diff --git a/flake.nix b/flake.nix
index f0b0d3d..dad1ba8 100644
--- a/flake.nix
+++ b/flake.nix
@@ -30,6 +30,8 @@
               cudaPackages.cudatoolkit
               cudaPackages.cuda_nvcc
             ];
+            CUDA_ROOT = "${pkgs.cudaPackages.cudatoolkit}";
+            LD_LIBRARY_PATH = "/run/opengl-driver/lib";
           };
 
         }