fix cuda gnet_expand

src/hvm.cu

@@ -1852,6 +1852,11 @@ Port gnet_vars_load(GNet* gnet, u32 loc) {
   return port;
 }
 
+// Writes a host var to device
+void gnet_vars_create(GNet* gnet, u32 var, Port val) {
+  cudaMemcpy(&gnet->vars_buf[var], &val, sizeof(Port), cudaMemcpyHostToDevice);
+}
+
 // Like the enter() function, but from host and read-only
 Port gnet_peek(GNet* gnet, Port port) {
   while (get_tag(port) == VAR) {
@@ -1864,13 +1869,15 @@ Port gnet_peek(GNet* gnet, Port port) {
 
 // Expands a REF Port.
 Port gnet_expand(GNet* gnet, Port port) {
+  Port old = gnet_vars_load(gnet, get_val(ROOT));
   Port got = gnet_peek(gnet, port);
   //printf("expand %s\n", show_port(got).x);
   while (get_tag(got) == REF) {
     gnet_boot_redex(gnet, new_pair(new_port(REF,get_val(got)), ROOT));
     gnet_normalize(gnet);
     got = gnet_peek(gnet, gnet_vars_load(gnet, get_val(ROOT)));
   }
+  gnet_vars_create(gnet, get_val(ROOT), old);
   return got;
 }
 

tests/programs/hello-world.hvm

@@ -0,0 +1,23 @@
+@String/Cons = (a (b ((@String/Cons/tag (a (b c))) c)))
+
+@String/Cons/tag = 1
+
+@String/Nil = ((@String/Nil/tag a) a)
+
+@String/Nil/tag = 0
+
+@main = l
+  & @String/Cons ~ (104 (k l))
+  & @String/Cons ~ (101 (j k))
+  & @String/Cons ~ (108 (i j))
+  & @String/Cons ~ (108 (h i))
+  & @String/Cons ~ (111 (g h))
+  & @String/Cons ~ (44 (f g))
+  & @String/Cons ~ (32 (e f))
+  & @String/Cons ~ (119 (d e))
+  & @String/Cons ~ (111 (c d))
+  & @String/Cons ~ (114 (b c))
+  & @String/Cons ~ (108 (a b))
+  & @String/Cons ~ (100 (@String/Nil a))
+
+

tests/programs/list.hvm

@@ -0,0 +1,30 @@
+@List/Cons = (a (b ((@List/Cons/tag (a (b c))) c)))
+
+@List/Cons/tag = 1
+
+@List/Nil = ((@List/Nil/tag a) a)
+
+@List/Nil/tag = 0
+
+@id = (a a)
+
+@list = c
+  & @List/Cons ~ (1 (b c))
+  & @List/Cons ~ (2 (a b))
+  & @List/Cons ~ (3 (@List/Nil a))
+
+@main = e
+  & @map ~ ((a b) (d e))
+  & @map ~ (a (@list b))
+  & @List/Cons ~ (@id (@List/Nil d))
+
+// @main__C0 = (a b)
+//   & @map ~ (a (@list b))
+
+@map = (a ((@map__C1 (a b)) b))
+
+@map__C0 = (* (a (d ({(a b) c} f))))
+  & @List/Cons ~ (b (e f))
+  & @map ~ (c (d e))
+
+@map__C1 = (?(((* @List/Nil) @map__C0) a) a)

tests/snapshots/run__file@hello-world.hvm.snap

@@ -0,0 +1,6 @@
+---
+source: tests/run.rs
+expression: rust_output
+input_file: tests/programs/hello-world.hvm
+---
+Result: ((@String/Cons/tag (104 (((@String/Cons/tag (101 (((@String/Cons/tag (108 (((@String/Cons/tag (108 (((@String/Cons/tag (111 (((@String/Cons/tag (44 (((@String/Cons/tag (32 (((@String/Cons/tag (119 (((@String/Cons/tag (111 (((@String/Cons/tag (114 (((@String/Cons/tag (108 (((@String/Cons/tag (100 (@String/Nil x0))) x0) x1))) x1) x2))) x2) x3))) x3) x4))) x4) x5))) x5) x6))) x6) x7))) x7) x8))) x8) x9))) x9) x10))) x10) x11))) x11)

tests/snapshots/run__file@list.hvm.snap

@@ -0,0 +1,6 @@
+---
+source: tests/run.rs
+expression: rust_output
+input_file: tests/programs/list.hvm
+---
+Result: ((@List/Cons/tag (((@List/Cons/tag (1 (((@List/Cons/tag (* (((@List/Cons/tag (* (@List/Nil x0))) x0) x1))) x1) x2))) x2) (@List/Nil x3))) x3)