avoid huge CUDA memset - faster boot

this is a hotfix for huge load times in some devices
HigherOrderCO · May 23, 2024 · 38fc2f1 · 38fc2f1
1 parent 656c881
commit 38fc2f1
Show file tree

Hide file tree

Showing 2 changed files with 16 additions and 2 deletions.
diff --git a/Cargo.toml b/Cargo.toml
@@ -1,6 +1,6 @@
 [package]
 name = "hvm"
-version = "2.0.14"
+version = "2.0.15"
 edition = "2021"
 build = "build.rs"
 description = "A massively parallel, optimal functional runtime in Rust."

diff --git a/src/hvm.cu b/src/hvm.cu
@@ -1685,10 +1685,24 @@ __global__ void evaluator(GNet* gnet) {
 // GNet Host Functions
 // -------------------
 
+// Initializes the GNet
+__global__ void initialize(GNet* gnet) {
+  gnet->node_put[GID()] = 0;
+  gnet->vars_put[GID()] = 0;
+  gnet->rbag_pos[GID()] = 0;
+  for (u32 i = 0; i < RLEN; ++i) {
+    gnet->rbag_buf_A[G_RBAG_LEN / TPG * GID()] = 0;
+  }
+  for (u32 i = 0; i < RLEN; ++i) {
+    gnet->rbag_buf_B[G_RBAG_LEN / TPG * GID()] = 0;
+  }
+}
+
 GNet* gnet_create() {
   GNet *gnet;
   cudaMalloc((void**)&gnet, sizeof(GNet));
-  cudaMemset(gnet, 0, sizeof(GNet));
+  initialize<<<BPG, TPB>>>(gnet);
+  //cudaMemset(gnet, 0, sizeof(GNet));
   return gnet;
 }