style: remove structs that have been made redundant

Author: 876pol

crates/vm/cuda/include/system/records.cuh

@@ -0,0 +1,11 @@
+#pragma once
+
+#include <cstdint>
+#include <cstddef>
+
+template <size_t CHUNK, size_t BLOCKS> struct MemoryInventoryRecord {
+    uint32_t address_space;
+    uint32_t ptr;
+    uint32_t timestamps[BLOCKS];
+    uint32_t values[CHUNK];
+};

crates/vm/cuda/src/system/boundary.cu

@@ -3,19 +3,13 @@
 #include "primitives/less_than.cuh"
 #include "primitives/shared_buffer.cuh"
 #include "primitives/trace_access.h"
+#include "system/records.cuh"
 #include <cassert>
 
 inline constexpr size_t PERSISTENT_CHUNK = 8;
 inline constexpr size_t BLOCKS_PER_CHUNK = 2;
 inline constexpr size_t VOLATILE_CHUNK = 1;
 
-template <size_t CHUNK, size_t BLOCKS> struct BoundaryRecord {
-    uint32_t address_space;
-    uint32_t ptr;
-    uint32_t timestamps[BLOCKS];
-    uint32_t values[CHUNK];
-};
-
 template <typename T> struct PersistentBoundaryCols {
     T expand_direction;
     T address_space;
@@ -43,7 +37,7 @@ __global__ void cukernel_persistent_boundary_tracegen(
     size_t height,
     size_t width,
     uint8_t const *const *initial_mem,
-    BoundaryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> *records,
+    MemoryInventoryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> *records,
     size_t num_records,
     FpArray<16> *poseidon2_buffer,
     uint32_t *poseidon2_buffer_idx,
@@ -54,7 +48,7 @@ __global__ void cukernel_persistent_boundary_tracegen(
     RowSlice row = RowSlice(trace + row_idx, height);
 
     if (record_idx < num_records) {
-        BoundaryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> record = records[record_idx];
+        MemoryInventoryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> record = records[record_idx];
         Poseidon2Buffer poseidon2(poseidon2_buffer, poseidon2_buffer_idx, poseidon2_capacity);
         COL_WRITE_VALUE(row, PersistentBoundaryCols, address_space, record.address_space);
         COL_WRITE_VALUE(row, PersistentBoundaryCols, leaf_label, record.ptr / PERSISTENT_CHUNK);
@@ -114,7 +108,7 @@ __global__ void cukernel_volatile_boundary_tracegen(
     Fp *trace,
     size_t height,
     size_t width,
-    BoundaryRecord<VOLATILE_CHUNK, 1> const *records,
+    MemoryInventoryRecord<VOLATILE_CHUNK, 1> const *records,
     size_t num_records,
     uint32_t *range_checker,
     size_t range_checker_num_bins,
@@ -131,7 +125,7 @@ __global__ void cukernel_volatile_boundary_tracegen(
             // For the sake of always filling `addr_lt_aux`
             row.fill_zero(0, width);
         }
-        BoundaryRecord<VOLATILE_CHUNK, 1> record = records[idx];
+        MemoryInventoryRecord<VOLATILE_CHUNK, 1> record = records[idx];
         rc.decompose(
             record.address_space,
             as_max_bits,
@@ -150,7 +144,7 @@ __global__ void cukernel_volatile_boundary_tracegen(
         COL_WRITE_VALUE(row, VolatileBoundaryCols, is_valid, Fp::one());
 
         if (idx != num_records - 1) {
-            BoundaryRecord<VOLATILE_CHUNK, 1> next_record = records[idx + 1];
+            MemoryInventoryRecord<VOLATILE_CHUNK, 1> next_record = records[idx + 1];
             uint32_t curr[ADDR_ELTS] = {record.address_space, record.ptr};
             uint32_t next[ADDR_ELTS] = {next_record.address_space, next_record.ptr};
             IsLessThanArray::generate_subrow(
@@ -198,8 +192,8 @@ extern "C" int _persistent_boundary_tracegen(
     size_t poseidon2_capacity
 ) {
     auto [grid, block] = kernel_launch_params(height);
-    BoundaryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> *d_records =
-        reinterpret_cast<BoundaryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> *>(d_raw_records);
+    MemoryInventoryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> *d_records =
+        reinterpret_cast<MemoryInventoryRecord<PERSISTENT_CHUNK, BLOCKS_PER_CHUNK> *>(d_raw_records);
     FpArray<16> *d_poseidon2_buffer = reinterpret_cast<FpArray<16> *>(d_poseidon2_raw_buffer);
     cukernel_persistent_boundary_tracegen<<<grid, block>>>(
         d_trace,
@@ -227,7 +221,8 @@ extern "C" int _volatile_boundary_tracegen(
     size_t ptr_max_bits
 ) {
     auto [grid, block] = kernel_launch_params(height, 512);
-    auto d_records = reinterpret_cast<BoundaryRecord<VOLATILE_CHUNK, 1> const *>(d_raw_records);
+    auto d_records =
+        reinterpret_cast<MemoryInventoryRecord<VOLATILE_CHUNK, 1> const *>(d_raw_records);
     cukernel_volatile_boundary_tracegen<<<grid, block>>>(
         d_trace,
         height,

crates/vm/cuda/src/system/inventory.cu

@@ -1,23 +1,15 @@
 #include "launcher.cuh"
+#include "system/records.cuh"
 #include <cub/device/device_scan.cuh>
 #include <cstddef>
 #include <cstdint>
 
-template <size_t CHUNK, size_t BLOCKS> struct MemoryInventoryRecord {
-    uint32_t address_space;
-    uint32_t ptr;
-    uint32_t timestamps[BLOCKS];
-    uint32_t values[CHUNK];
-};
-
-
 const uint32_t IN_BLOCK_SIZE = 4;
 const uint32_t OUT_BLOCK_SIZE = 8;
 
 using InRec = MemoryInventoryRecord<IN_BLOCK_SIZE, 1>;
 using OutRec = MemoryInventoryRecord<OUT_BLOCK_SIZE, 2>;
 
-
 __device__ inline bool same_output_block(
     InRec const *in,
     size_t lhs_idx,

crates/vm/cuda/src/system/memory/merkle_tree.cu

@@ -2,6 +2,7 @@
 #include "poseidon2.cuh"
 #include "primitives/shared_buffer.cuh"
 #include "primitives/trace_access.h"
+#include "system/records.cuh"
 
 #include <cub/cub.cuh>
 
@@ -159,12 +160,7 @@ template <typename T> struct MerkleCols {
     T right_direction_different;
 };
 
-struct LabeledDigest {
-    uint32_t address_space;
-    uint32_t ptr;
-    uint32_t timestamps[2];
-    uint32_t digest_raw[CELLS_OUT];
-};
+using LabeledDigest = MemoryInventoryRecord<CELLS_OUT, 2>;
 
 __global__ void prepare_for_updating(
     uint32_t *child_buf,
@@ -177,9 +173,9 @@ __global__ void prepare_for_updating(
     }
     child_buf[gid] = gid;
     Fp cells[CELLS] = {0};
-    COPY_DIGEST(cells, leaves[gid].digest_raw);
+    COPY_DIGEST(cells, leaves[gid].values);
     poseidon2_mix(cells);
-    COPY_DIGEST(leaves[gid].digest_raw, cells);
+    COPY_DIGEST(leaves[gid].values, cells);
     leaves[gid].address_space -= 1;
     leaves[gid].ptr /= CELLS_OUT;
     leaves[gid].timestamps[0] =
@@ -339,13 +335,13 @@ __global__ void update_merkle_layer(
     { // new values trace row + actual update
         bool left_new = false;
         if (auto const child_ptr = child_ptrs[2 * idx]; child_ptr != MISSING_CHILD) {
-            COPY_DIGEST(&subtree_layer[2 * parent_label], layer[child_ptr].digest_raw);
+            COPY_DIGEST(&subtree_layer[2 * parent_label], layer[child_ptr].values);
             left_new = true;
         }
         COPY_DIGEST(cells, old_left_digest);
         bool right_new = false;
         if (auto const child_ptr = child_ptrs[2 * idx + 1]; child_ptr != MISSING_CHILD) {
-            COPY_DIGEST(&subtree_layer[2 * parent_label + 1], layer[child_ptr].digest_raw);
+            COPY_DIGEST(&subtree_layer[2 * parent_label + 1], layer[child_ptr].values);
             right_new = true;
         }
         COPY_DIGEST(cells + CELLS_OUT, old_right_digest);
@@ -361,7 +357,7 @@ __global__ void update_merkle_layer(
             right_new,
             poseidon2
         );
-        COPY_DIGEST(layer[parent_ptr].digest_raw, cells);
+        COPY_DIGEST(layer[parent_ptr].values, cells);
     }
 }
 
@@ -392,7 +388,7 @@ __global__ void update_to_root(
         auto const address_space_idx = layer[idx].address_space;
         layer[idx].ptr = num_roots - 1 + address_space_idx;
         if (subtrees[address_space_idx]) {
-            COPY_DIGEST(subtrees[address_space_idx], layer[idx].digest_raw);
+            COPY_DIGEST(subtrees[address_space_idx], layer[idx].values);
         }
     }
 
@@ -425,7 +421,7 @@ __global__ void update_to_root(
         for (auto i : {0, 1}) {
             if (children_ids[i] != MISSING_CHILD) {
                 COPY_DIGEST(
-                    &out[2 * out_idx + 1 + i], layer[layer_ids[children_ids[i]]].digest_raw
+                    &out[2 * out_idx + 1 + i], layer[layer_ids[children_ids[i]]].values
                 );
             }
             COPY_DIGEST(cells + CELLS_OUT * i, &out[2 * out_idx + 1 + i]);
@@ -450,11 +446,11 @@ __global__ void update_to_root(
                 children_ids[1] != MISSING_CHILD,
                 poseidon2
             );
-            COPY_DIGEST(layer[layer_ids[surely_surviving_child]].digest_raw, cells);
+            COPY_DIGEST(layer[layer_ids[surely_surviving_child]].values, cells);
             COL_WRITE_VALUE(row, MerkleCols, height_section, true);
         }
     }
-    COPY_DIGEST(out, layer[layer_ids[0]].digest_raw);
+    COPY_DIGEST(out, layer[layer_ids[0]].values);
     for (auto i : {0, 1}) {
         RowSlice row(merkle_trace + i, trace_height);
         COL_WRITE_VALUE(row, MerkleCols, is_root, true);