From 7c6aeb76e914efbb4362e6aa992fc23dd1e111ba Mon Sep 17 00:00:00 2001
From: prrathi <prrathi10@gmail.com>
Date: Sat, 18 Jan 2025 16:05:47 -0600
Subject: [PATCH] extra
---
.gitignore | 3 +-
Cargo.lock | 8 ++--
Cargo.toml | 4 +-
hercules_cg/src/gpu.rs | 44 ++++++++++++++-----
hercules_ir/src/dot.rs | 3 +-
hercules_opt/src/pass.rs | 6 +++
hercules_samples/call/Cargo.toml | 7 ++-
hercules_samples/ccp/build.rs | 7 ---
hercules_samples/dot/build.rs | 7 ---
hercules_samples/fac/build.rs | 7 ---
hercules_samples/matmul/build.rs | 7 ---
hercules_samples/matmul/src/matmul.hir | 10 ++---
juno_build/build.rs | 8 ++++
juno_build/src/lib.rs | 2 +-
juno_samples/antideps/Cargo.toml | 9 ++--
juno_samples/antideps/antideps.mod | Bin 0 -> 1583 bytes
juno_samples/casts_and_intrinsics/Cargo.toml | 7 ++-
juno_samples/cava/Cargo.toml | 6 +--
juno_samples/concat/Cargo.toml | 6 +--
juno_samples/implicit_clone/Cargo.toml | 9 ++--
juno_samples/matmul/Cargo.toml | 9 ++--
juno_samples/nested_ccp/Cargo.toml | 9 ++--
juno_samples/simple3/Cargo.toml | 9 ++--
23 files changed, 109 insertions(+), 78 deletions(-)
create mode 100644 juno_build/build.rs
create mode 100644 juno_samples/antideps/antideps.mod
diff --git a/.gitignore b/.gitignore
index 45f2e61b..29eb3e04 100644
--- a/.gitignore
+++ b/.gitignore
@@ -4,11 +4,10 @@
*.out
*.ll
*.c
+*.cu
*.o
*.a
*.hrt
.*.swp
.vscode
*_env
-
-juno_samples/matmul/src/matmul_indented.jn
diff --git a/Cargo.lock b/Cargo.lock
index a1eb77de..3f1e2e7c 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -374,9 +374,9 @@ dependencies = [
[[package]]
name = "cc"
-version = "1.2.9"
+version = "1.2.10"
source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "c8293772165d9345bdaaa39b45b2109591e63fe5e6fbc23c6ff930a048aa310b"
+checksum = "13208fcbb66eaeffe09b99fffbe1af420f00a7b35aa99ad683dfc1aa76145229"
dependencies = [
"jobserver",
"libc",
@@ -1278,9 +1278,9 @@ checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a"
[[package]]
name = "miniz_oxide"
-version = "0.8.2"
+version = "0.8.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "4ffbe83022cedc1d264172192511ae958937694cd57ce297164951b8b3568394"
+checksum = "b8402cab7aefae129c6977bb0ff1b8fd9a04eb5b51efc50a70bea51cda0c7924"
dependencies = [
"adler2",
"simd-adler32",
diff --git a/Cargo.toml b/Cargo.toml
index c57125f7..4e5826ca 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -5,13 +5,13 @@ members = [
"hercules_ir",
"hercules_opt",
"hercules_rt",
-
+
"hercules_tools/hercules_driver",
"juno_frontend",
"juno_scheduler",
"juno_build",
-
+
#"hercules_test/hercules_interpreter",
#"hercules_test/hercules_tests",
diff --git a/hercules_cg/src/gpu.rs b/hercules_cg/src/gpu.rs
index 731e57e3..31c50212 100644
--- a/hercules_cg/src/gpu.rs
+++ b/hercules_cg/src/gpu.rs
@@ -638,7 +638,10 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
let ret_type_pnt = self.get_type(*self.return_type_id, true);
write!(w, "\t{} ret;\n", ret_type_pnt)?;
write!(w, "\tcudaMalloc((void**)&ret, sizeof({}));\n", ret_type)?;
- write!(pass_args, ", ret")?;
+ if !first_param {
+ write!(pass_args, ", ")?;
+ }
+ write!(pass_args, "ret")?;
}
write!(w, "\t{}_gpu<<<{}, {}, {}>>>({});\n", self.function.name, num_blocks, num_threads, dynamic_shared_offset, pass_args)?;
write!(w, "\tcudaDeviceSynchronize();\n")?;
@@ -897,6 +900,9 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
for data in self.bbs.1[control.idx()].iter() {
self.codegen_data_node(*data, KernelState::OutBlock, None, None, None, false, extra_dim_collects, dynamic_shared_offset, body, &mut tabs)?;
}
+ for data in self.bbs.1[control.idx()].iter() {
+ self.codegen_data_phi(*data, tabs, body)?;
+ }
Ok(())
})
}
@@ -929,6 +935,9 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
for data in self.bbs.1[control.idx()].iter() {
self.codegen_data_node(*data, state, None, None, None, false, extra_dim_collects, dynamic_shared_offset, body, &mut tabs)?;
}
+ for data in self.bbs.1[control.idx()].iter() {
+ self.codegen_data_phi(*data, tabs, body)?;
+ }
}
// Then generate data and control for the single block fork if it exists
if block_fork.is_some() {
@@ -943,6 +952,9 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
for data in self.bbs.1[control.idx()].iter() {
self.codegen_data_node(*data, state, Some(num_threads), None, Some(block_fork.unwrap()), false, extra_dim_collects, dynamic_shared_offset, body, &mut tabs)?;
}
+ for data in self.bbs.1[control.idx()].iter() {
+ self.codegen_data_phi(*data, tabs, body)?;
+ }
}
}
// Then generate for the thread fork tree through Fork node traversal.
@@ -1022,6 +1034,9 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
&mut tabs,
)?;
}
+ for data in self.bbs.1[control.idx()].iter() {
+ self.codegen_data_phi(*data, tabs, body)?;
+ }
}
for child in fork_tree.get(&curr_fork).unwrap() {
self.codegen_data_control_traverse(
@@ -1398,15 +1413,8 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
panic!("Unsupported data node type")
}
}
- // Since the data uses and reducts are responsible for updating Phi and
- // Reduce nodes, respectively, we check and emit those for each data node.
- if let Some(phis) = self.label_data_for_phi.get(&id) {
- let val = self.get_value(id, false, false);
- for phi in phis {
- let phi_val = self.get_value(*phi, false, false);
- write!(w, "{}{} = {};\n", tabs, phi_val, val,)?;
- }
- }
+ // Since reducts are responsible for updating Reduce nodes,
+ // we check and emit those for each data node.
if let Some(reduces) = self.reduct_reduce_map.get(&id) {
let val = self.get_value(id, false, false);
for reduce in reduces {
@@ -1417,6 +1425,22 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
Ok(())
}
+ /*
+ * Update Phi assignments for each data node. This is run after all data nodes
+ * for given control block have been emitted.
+ */
+ fn codegen_data_phi(&self, id: NodeID, num_tabs: usize, w: &mut String) -> Result<(), Error> {
+ let tabs = "\t".repeat(num_tabs);
+ if let Some(phis) = self.label_data_for_phi.get(&id) {
+ let val = self.get_value(id, false, false);
+ for phi in phis {
+ let phi_val = self.get_value(*phi, false, false);
+ write!(w, "{}{} = {};\n", tabs, phi_val, val)?;
+ }
+ }
+ Ok(())
+ }
+
fn codegen_control_node(
&self,
id: NodeID,
diff --git a/hercules_ir/src/dot.rs b/hercules_ir/src/dot.rs
index 4d526366..fe6fee09 100644
--- a/hercules_ir/src/dot.rs
+++ b/hercules_ir/src/dot.rs
@@ -1,5 +1,5 @@
use std::collections::HashMap;
-use std::env::temp_dir;
+use std::env::{temp_dir};
use std::fmt::Write;
use std::fs::File;
use std::io::Write as _;
@@ -23,6 +23,7 @@ pub fn xdot_module(
let mut rng = rand::thread_rng();
let num: u64 = rng.gen();
tmp_path.push(format!("hercules_dot_{}.dot", num));
+ let tmp_path = std::path::PathBuf::from(format!("hercules_dot.dot"));
let mut file = File::create(&tmp_path).expect("PANIC: Unable to open output file.");
let mut contents = String::new();
write_dot(
diff --git a/hercules_opt/src/pass.rs b/hercules_opt/src/pass.rs
index 149e4eeb..dbc24016 100644
--- a/hercules_opt/src/pass.rs
+++ b/hercules_opt/src/pass.rs
@@ -1079,6 +1079,12 @@ impl PassManager {
file.write_all(cuda_ir.as_bytes())
.expect("PANIC: Unable to write output CUDA IR file contents.");
+ let cuda_text_path = format!("{}.cu", module_name);
+ let mut cuda_text_file = File::create(&cuda_text_path)
+ .expect("PANIC: Unable to open CUDA IR text file.");
+ cuda_text_file.write_all(cuda_ir.as_bytes())
+ .expect("PANIC: Unable to write CUDA IR text file contents.");
+
let mut nvcc_process = Command::new("nvcc")
.arg("-c")
.arg("-O3")
diff --git a/hercules_samples/call/Cargo.toml b/hercules_samples/call/Cargo.toml
index 52c588e7..a5a44c2e 100644
--- a/hercules_samples/call/Cargo.toml
+++ b/hercules_samples/call/Cargo.toml
@@ -4,11 +4,14 @@ version = "0.1.0"
authors = ["Russel Arbore <rarbore2@illinois.edu>"]
edition = "2021"
+[features]
+cuda = ["juno_build/cuda"]
+
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
rand = "*"
async-std = "*"
with_builtin_macros = "0.1.0"
diff --git a/hercules_samples/ccp/build.rs b/hercules_samples/ccp/build.rs
index 0b984a0f..f04d48c7 100644
--- a/hercules_samples/ccp/build.rs
+++ b/hercules_samples/ccp/build.rs
@@ -6,11 +6,4 @@ fn main() {
.unwrap()
.build()
.unwrap();
-
- #[cfg(feature = "cuda")]
- println!("cargo::rustc-link-search=native=/usr/lib/x86_64-linux-gnu/");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-search=native=/usr/local/cuda/lib64");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-lib=cudart");
}
diff --git a/hercules_samples/dot/build.rs b/hercules_samples/dot/build.rs
index 43cd34f9..2a239bc6 100644
--- a/hercules_samples/dot/build.rs
+++ b/hercules_samples/dot/build.rs
@@ -6,11 +6,4 @@ fn main() {
.unwrap()
.build()
.unwrap();
-
- #[cfg(feature = "cuda")]
- println!("cargo::rustc-link-search=native=/usr/lib/x86_64-linux-gnu/");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-search=native=/usr/local/cuda/lib64");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-lib=cudart");
}
diff --git a/hercules_samples/fac/build.rs b/hercules_samples/fac/build.rs
index 6863b27b..4d8226f1 100644
--- a/hercules_samples/fac/build.rs
+++ b/hercules_samples/fac/build.rs
@@ -6,11 +6,4 @@ fn main() {
.unwrap()
.build()
.unwrap();
-
- #[cfg(feature = "cuda")]
- println!("cargo::rustc-link-search=native=/usr/lib/x86_64-linux-gnu/");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-search=native=/usr/local/cuda/lib64");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-lib=cudart");
}
diff --git a/hercules_samples/matmul/build.rs b/hercules_samples/matmul/build.rs
index b170024b..08478dea 100644
--- a/hercules_samples/matmul/build.rs
+++ b/hercules_samples/matmul/build.rs
@@ -6,11 +6,4 @@ fn main() {
.unwrap()
.build()
.unwrap();
-
- #[cfg(feature = "cuda")]
- println!("cargo::rustc-link-search=native=/usr/lib/x86_64-linux-gnu/");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-search=native=/usr/local/cuda/lib64");
- #[cfg(feature = "cuda")]
- println!("cargo:rustc-link-lib=cudart");
}
diff --git a/hercules_samples/matmul/src/matmul.hir b/hercules_samples/matmul/src/matmul.hir
index 400ab5e1..b0c31da4 100644
--- a/hercules_samples/matmul/src/matmul.hir
+++ b/hercules_samples/matmul/src/matmul.hir
@@ -1,9 +1,9 @@
-fn matmul(a: array(i32, 16, 64), b: array(i32, 64, 32)) -> array(i32, 16, 32)
- c = constant(array(i32, 16, 32), [])
- i_j_ctrl = fork(start, 16, 32)
+fn matmul<3>(a: array(i32, #0, #1), b: array(i32, #1, #2)) -> array(i32, #0, #2)
+ c = constant(array(i32, #0, #2), [])
+ i_j_ctrl = fork(start, #0, #2)
i_idx = thread_id(i_j_ctrl, 0)
j_idx = thread_id(i_j_ctrl, 1)
- k_ctrl = fork(i_j_ctrl, 64)
+ k_ctrl = fork(i_j_ctrl, #1)
k_idx = thread_id(k_ctrl, 0)
k_join_ctrl = join(k_ctrl)
i_j_join_ctrl = join(k_join_ctrl)
@@ -15,4 +15,4 @@ fn matmul(a: array(i32, 16, 64), b: array(i32, 64, 32)) -> array(i32, 16, 32)
add = add(mul, dot)
dot = reduce(k_join_ctrl, zero, add)
update_c = write(update_i_j_c, dot, position(i_idx, j_idx))
- update_i_j_c = reduce(i_j_join_ctrl, c, update_c)
\ No newline at end of file
+ update_i_j_c = reduce(i_j_join_ctrl, c, update_c)
diff --git a/juno_build/build.rs b/juno_build/build.rs
new file mode 100644
index 00000000..7ba34c8c
--- /dev/null
+++ b/juno_build/build.rs
@@ -0,0 +1,8 @@
+fn main() {
+ #[cfg(feature = "cuda")]
+ println!("cargo::rustc-link-search=native=/usr/lib/x86_64-linux-gnu/");
+ #[cfg(feature = "cuda")]
+ println!("cargo:rustc-link-search=native=/usr/local/cuda/lib64");
+ #[cfg(feature = "cuda")]
+ println!("cargo:rustc-link-lib=cudart");
+}
diff --git a/juno_build/src/lib.rs b/juno_build/src/lib.rs
index 0c676e4c..40660806 100644
--- a/juno_build/src/lib.rs
+++ b/juno_build/src/lib.rs
@@ -27,7 +27,7 @@ impl JunoCompiler {
src_path: None,
out_path: None,
verify: JunoVerify::None,
- x_dot: false,
+ x_dot: true,
schedule: JunoSchedule::None,
}
}
diff --git a/juno_samples/antideps/Cargo.toml b/juno_samples/antideps/Cargo.toml
index e492e2ae..e6f38e09 100644
--- a/juno_samples/antideps/Cargo.toml
+++ b/juno_samples/antideps/Cargo.toml
@@ -8,11 +8,14 @@ edition = "2021"
name = "juno_antideps"
path = "src/main.rs"
+[features]
+cuda = ["juno_build/cuda", "hercules_rt/cuda"]
+
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
-hercules_rt = { path = "../../hercules_rt", features = [] }
+juno_build = { path = "../../juno_build" }
+hercules_rt = { path = "../../hercules_rt" }
with_builtin_macros = "0.1.0"
async-std = "*"
diff --git a/juno_samples/antideps/antideps.mod b/juno_samples/antideps/antideps.mod
new file mode 100644
index 0000000000000000000000000000000000000000..b4abaef48222ace6f66264d6746dd2d7924b9039
GIT binary patch
literal 1583
zcmb7E{c;;M5Z_(h>2#-)e3s;|IEihN*s;?z{z#gH4$wF8K$|~0Q%Hs~3_KnT!vN)f
zUV*oO)nP*8!c4gt?d`7iOFyl)XuN*;`Q59V^`~Ea_438d-D?6szzhl~Fi#{RXsD7D
z1Yw%N<}<HuZ}0A1kd08XM=sK6kCl)>n8=NF^Z9K|@y)%%qV|Y6a#3f{x~xP70+WAo
zyyLap&eN`hdOJ_+>zgmXU2i8tt^!q^D~r{2bx-#zkwe1Sr26|v(A)xIJr@8ed)pcd
zXOIvrVeXA$mRwpZ_^6E}FF*^OaQIwj1=;(ce=&^*aIF`FlH=}jT;FeUlwi?k1R^9D
z60nwqUKSJv1+G=^YX`nn)@<9O>(pGc0b1!o=bB`X$S^@GyUZcU0x8Y&ZdKs!Ry7Yr
zzm3;oy%cCH7kJGDx+_k#Z}$lNrC@eaD<&3)T(Aroh74?<g^MC+3>y6I2W&R7vxAJz
z8==Bbfjfzwb~XhQ2%Kf9WqljGJ!1CIpN!+anGvuJ`uT#zrA{<ys~ft~O$OI#uY)$a
zr3-mZ*k_WQ$_uhxCMP|)qV$Zp&X_l@9x)jlGnwol<qU^n?3ylLJua}I>xT&ZAL$*m
zmg^N^gTb1{+dSrUIPlcx-=2$j4Q$c}i@R^3h5H$)4O+It!jTOP3=MF@=w{sr{Lt6F
zp4kvPV9$>B(M4njHtoi1?`X+r#H&9d(Pn}jIi>?44_QPq+*f6VbeQMkUV+p1KV;`$
zU(xzJfvV{_%ricy#p9HKg1{ia8WyfLz2ML0T;$ww`yu*;j)I-fz%8S~+M}GVH(lXB
znX&HF#iO2jKjmMj+kNI%%%9kH2khWcMhCsHgJB2!`Goade^LMJlPA3WIHRY-Fv2hb
zJ7Hjk+$bY%nFGD6kGX-2p2{ows9=JI$;YL<jN)bFp3-WMNo$|Ucn}{9Qo@h>(jmn*
zZuavJ4`as^;@SHH_#c+tJSm~%!BH}DmPZ9IB$3KYF&8z0S7HHEUXoYFOCgkqCBmsm
z4Nogk2&;Ki@Vp_JibMoK#u%$m{D`7dh4=xL7T=?&#ExkQX)4n+&Ulq5qe>1a)N}Qj
tnyELa{`l^<GxaN~-!qk|7djzIbsMGT@se{2gak?;c!Tf@!p{gl;h$>7iFE(~
literal 0
HcmV?d00001
diff --git a/juno_samples/casts_and_intrinsics/Cargo.toml b/juno_samples/casts_and_intrinsics/Cargo.toml
index 83d5be58..9fac18b7 100644
--- a/juno_samples/casts_and_intrinsics/Cargo.toml
+++ b/juno_samples/casts_and_intrinsics/Cargo.toml
@@ -8,10 +8,13 @@ edition = "2021"
name = "juno_casts_and_intrinsics"
path = "src/main.rs"
+[features]
+cuda = ["juno_build/cuda"]
+
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
with_builtin_macros = "0.1.0"
async-std = "*"
diff --git a/juno_samples/cava/Cargo.toml b/juno_samples/cava/Cargo.toml
index dfde5978..ff375d80 100644
--- a/juno_samples/cava/Cargo.toml
+++ b/juno_samples/cava/Cargo.toml
@@ -9,11 +9,11 @@ name = "juno_cava"
path = "src/main.rs"
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
-hercules_rt = { path = "../../hercules_rt", features = [] }
+juno_build = { path = "../../juno_build" }
+hercules_rt = { path = "../../hercules_rt" }
async-std = "*"
clap = { version = "*", features = ["derive"] }
image = "*"
diff --git a/juno_samples/concat/Cargo.toml b/juno_samples/concat/Cargo.toml
index 888a083f..24ba1acf 100644
--- a/juno_samples/concat/Cargo.toml
+++ b/juno_samples/concat/Cargo.toml
@@ -9,10 +9,10 @@ name = "juno_concat"
path = "src/main.rs"
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
-hercules_rt = { path = "../../hercules_rt", features = [] }
+juno_build = { path = "../../juno_build" }
+hercules_rt = { path = "../../hercules_rt" }
with_builtin_macros = "0.1.0"
async-std = "*"
diff --git a/juno_samples/implicit_clone/Cargo.toml b/juno_samples/implicit_clone/Cargo.toml
index 4f5387e7..96129371 100644
--- a/juno_samples/implicit_clone/Cargo.toml
+++ b/juno_samples/implicit_clone/Cargo.toml
@@ -8,11 +8,14 @@ edition = "2021"
name = "juno_implicit_clone"
path = "src/main.rs"
+[features]
+cuda = ["juno_build/cuda", "hercules_rt/cuda"]
+
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
-hercules_rt = { path = "../../hercules_rt", features = [] }
+juno_build = { path = "../../juno_build" }
+hercules_rt = { path = "../../hercules_rt" }
with_builtin_macros = "0.1.0"
async-std = "*"
diff --git a/juno_samples/matmul/Cargo.toml b/juno_samples/matmul/Cargo.toml
index 1e004dcd..eac83d15 100644
--- a/juno_samples/matmul/Cargo.toml
+++ b/juno_samples/matmul/Cargo.toml
@@ -8,12 +8,15 @@ edition = "2021"
name = "juno_matmul"
path = "src/main.rs"
+[features]
+cuda = ["juno_build/cuda", "hercules_rt/cuda"]
+
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
-hercules_rt = { path = "../../hercules_rt", features = [] }
+juno_build = { path = "../../juno_build" }
+hercules_rt = { path = "../../hercules_rt" }
with_builtin_macros = "0.1.0"
async-std = "*"
rand = "*"
diff --git a/juno_samples/nested_ccp/Cargo.toml b/juno_samples/nested_ccp/Cargo.toml
index bcf1fff8..5ee3f747 100644
--- a/juno_samples/nested_ccp/Cargo.toml
+++ b/juno_samples/nested_ccp/Cargo.toml
@@ -8,11 +8,14 @@ edition = "2021"
name = "juno_nested_ccp"
path = "src/main.rs"
+[features]
+cuda = ["juno_build/cuda", "hercules_rt/cuda"]
+
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
-hercules_rt = { path = "../../hercules_rt", features = [] }
+juno_build = { path = "../../juno_build" }
+hercules_rt = { path = "../../hercules_rt" }
with_builtin_macros = "0.1.0"
async-std = "*"
diff --git a/juno_samples/simple3/Cargo.toml b/juno_samples/simple3/Cargo.toml
index c66dc977..36d50dbd 100644
--- a/juno_samples/simple3/Cargo.toml
+++ b/juno_samples/simple3/Cargo.toml
@@ -8,11 +8,14 @@ edition = "2021"
name = "juno_simple3"
path = "src/main.rs"
+[features]
+cuda = ["juno_build/cuda", "hercules_rt/cuda"]
+
[build-dependencies]
-juno_build = { path = "../../juno_build", features = [] }
+juno_build = { path = "../../juno_build" }
[dependencies]
-juno_build = { path = "../../juno_build", features = [] }
-hercules_rt = { path = "../../hercules_rt", features = [] }
+juno_build = { path = "../../juno_build" }
+hercules_rt = { path = "../../hercules_rt" }
with_builtin_macros = "0.1.0"
async-std = "*"
--
GitLab