From 38fc2f1521c679154ac5c763b7bdabf756aa2cbf Mon Sep 17 00:00:00 2001
From: Victor Taelin <victor.taelin@gmail.com>
Date: Wed, 22 May 2024 21:51:50 -0300
Subject: [PATCH] avoid huge CUDA memset - faster boot

this is a hotfix for huge load times in some devices
---
 Cargo.toml |  2 +-
 src/hvm.cu | 16 +++++++++++++++-
 2 files changed, 16 insertions(+), 2 deletions(-)

diff --git a/Cargo.toml b/Cargo.toml
index bc4375d9..56a29f96 100644
--- 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
index 59510df7..7ade71a3 100644
--- 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;
 }