From 1d9c4edb0b4fa9694e7708858a5d856709e0a855 Mon Sep 17 00:00:00 2001
From: benlaflen <114254189+benlaflen@users.noreply.github.com>
Date: Thu, 8 Oct 2026 12:20:51 -0400
Subject: [PATCH 1/7] Unify mbarrier address decoding

---
 gpu-simulator/trace-parser/trace_parser.cc    | 83 +---------------
 gpu-simulator/trace-parser/trace_parser.h     | 18 +---
 util/tracer_nvbit/tracer_tool/decode.h        | 99 +++++++++++++++++++
 util/tracer_nvbit/tracer_tool/tracer_tool.cu  |  9 +-
 .../post-traces-processing.cpp                | 41 +-------
 5 files changed, 108 insertions(+), 142 deletions(-)
 create mode 100644 util/tracer_nvbit/tracer_tool/decode.h

diff --git a/gpu-simulator/trace-parser/trace_parser.cc b/gpu-simulator/trace-parser/trace_parser.cc
index 316da183f7..53a7905aa2 100644
--- a/gpu-simulator/trace-parser/trace_parser.cc
+++ b/gpu-simulator/trace-parser/trace_parser.cc
@@ -97,47 +97,6 @@ kernel_trace_t::kernel_trace_t(const std::string &filePath) {
 
 kernel_trace_t::~kernel_trace_t() { delete tracez_reader; }
 
-void inst_memadd_info_t::base_stride_decompress(
-    unsigned long long base_address, int stride,
-    const std::bitset<WARP_SIZE> &mask) {
-  bool first_bit1_found = false;
-  bool last_bit1_found = false;
-  unsigned long long addra = base_address;
-  for (int s = 0; s < WARP_SIZE; s++) {
-    if (mask.test(s) && !first_bit1_found) {
-      first_bit1_found = true;
-      addrs[s] = base_address;
-    } else if (first_bit1_found && !last_bit1_found) {
-      if (mask.test(s)) {
-        addra += stride;
-        addrs[s] = addra;
-      } else
-        last_bit1_found = true;
-    } else
-      addrs[s] = 0;
-  }
-}
-
-void inst_memadd_info_t::base_delta_decompress(
-    unsigned long long base_address, const std::vector<long long> &deltas,
-    const std::bitset<WARP_SIZE> &mask) {
-  bool first_bit1_found = false;
-  long long last_address = 0;
-  unsigned delta_index = 0;
-  for (int s = 0; s < 32; s++) {
-    if (mask.test(s) && !first_bit1_found) {
-      addrs[s] = base_address;
-      first_bit1_found = true;
-      last_address = base_address;
-    } else if (mask.test(s) && first_bit1_found) {
-      assert(delta_index < deltas.size());
-      addrs[s] = last_address + deltas[delta_index++];
-      last_address = addrs[s];
-    } else
-      addrs[s] = 0;
-  }
-}
-
 void tma_inst_memaddr_info_t::base_delta_decompress(
     unsigned long long base_address, const std::vector<long long> &deltas,
     const std::bitset<WARP_SIZE> &mask) {
@@ -284,46 +243,8 @@ bool inst_trace_t::parse_from_string(std::string trace, unsigned trace_version,
     memadd_info->width = get_datawidth_from_opcode(opcode_tokens);
 
     ss >> std::dec >> address_mode;
-    if (address_mode == address_format::list_all) {
-      // read addresses one by one from the file
-      for (int s = 0; s < WARP_SIZE; s++) {
-        if (mask_bits.test(s))
-          ss >> std::hex >> memadd_info->addrs[s];
-        else
-          memadd_info->addrs[s] = 0;
-      }
-    } else if (address_mode == address_format::base_stride) {
-      // read addresses as base address and stride
-      unsigned long long base_address = 0;
-      int stride = 0;
-      ss >> std::hex >> base_address;
-      ss >> std::dec >> stride;
-      memadd_info->base_stride_decompress(base_address, stride, mask_bits);
-    } else if (address_mode == address_format::base_delta) {
-      unsigned long long base_address = 0;
-      std::vector<long long> deltas;
-      // read addresses as base address and deltas
-      ss >> std::hex >> base_address;
-      // No delta will be pushed if there is only one bit set in the mask
-      // Find the first bit set
-      int first_bit_set = -1;
-      for (int s = 0; s < WARP_SIZE; s++) {
-        if (mask_bits.test(s)) {
-          first_bit_set = s;
-          break;
-        }
-      }
-      for (int s = first_bit_set + 1; s < WARP_SIZE; s++) {
-        if (mask_bits.test(s)) {
-          long long delta = 0;
-          ss >> std::dec >> delta;
-          deltas.push_back(delta);
-        }
-      }
-      memadd_info->base_delta_decompress(base_address, deltas, mask_bits);
-    } else {
-      assert(0 && "Unsupported address mode");
-    }
+    std::vector<uint64_t> addrs = decode(address_mode, ss, mask_bits); //TODO: JUST MAKE SURE THIS GETS IMPORTED
+    std::memcpy(memadd_info->addrs, addrs.data(), WARP_SIZE * sizeof(uint64_t));
   }
 
   ss >> imm;
diff --git a/gpu-simulator/trace-parser/trace_parser.h b/gpu-simulator/trace-parser/trace_parser.h
index 7e760644d5..8627052b1c 100644
--- a/gpu-simulator/trace-parser/trace_parser.h
+++ b/gpu-simulator/trace-parser/trace_parser.h
@@ -13,7 +13,8 @@
 #include <string>
 #include <vector>
 
-#define WARP_SIZE 32
+#include "decode.h"
+
 #define MAX_DST 1
 #define MAX_SRC 4
 
@@ -31,15 +32,6 @@ enum address_scope {
   SYS_MEM,
 };
 
-// TODO Fix this shared enum with tracer tool
-enum address_format {
-  list_all = 0,
-  base_stride = 1,
-  base_delta = 2,
-  tma_list_all = 3,
-  tma_base_delta = 4
-};
-
 struct trace_command {
   std::string command_string;
   command_type m_type;
@@ -48,12 +40,6 @@ struct trace_command {
 struct inst_memadd_info_t {
   uint64_t addrs[WARP_SIZE];
   int32_t width;
-
-  void base_stride_decompress(unsigned long long base_address, int stride,
-                              const std::bitset<WARP_SIZE> &mask);
-  void base_delta_decompress(unsigned long long base_address,
-                             const std::vector<long long> &deltas,
-                             const std::bitset<WARP_SIZE> &mask);
 };
 
 struct tma_inst_memaddr_info_t {
diff --git a/util/tracer_nvbit/tracer_tool/decode.h b/util/tracer_nvbit/tracer_tool/decode.h
new file mode 100644
index 0000000000..d22a5331b0
--- /dev/null
+++ b/util/tracer_nvbit/tracer_tool/decode.h
@@ -0,0 +1,99 @@
+#include <vector>
+#include <bitset>
+
+#define WARP_SIZE 32
+
+
+enum address_format {
+  list_all = 0,
+  base_stride = 1,
+  base_delta = 2,
+  tma_list_all = 3,
+  tma_base_delta = 4
+};
+
+inline vector<uint64_t> decode(
+    address_format address_mode,
+    std::stringstream &ss,
+    std::bitset<WARP_SIZE> mask_bits
+) {
+  switch(address_mode) {
+    case address_format::list_all: return decode_list_all(ss, mask_bits);
+    case address_format::base_stride: return decode_base_stride(ss, mask_bits);
+    case address_format::base_delta: return decode_base_delta(ss, mask_bits);
+    default: throw std::runtime_error("Unsupported address mode");
+  }
+}
+
+inline vector<uint64_t> decode_list_all(
+  std::stringstream &ss,
+  std::bitset<WARP_SIZE> mask_bits,
+) {
+  vector<uint64_t> addrs;
+  uint64_t addr = 0;
+
+  // read addresses one by one from the file
+  for (int s = 0; s < WARP_SIZE; s++) {
+    if (mask_bits.test(s)) {
+      ss >> std::hex >> addr;
+      addrs.push_back(addr);
+    } else addrs.push_back(0);
+  }
+
+  return addrs;
+}
+
+inline vector<uint64_t> decode_base_stride(
+  std::stringstream &ss,
+  std::bitset<WARP_SIZE> mask_bits
+) {
+  vector<uint64_t> addrs;
+  uint64_t addr = 0;
+
+  // read addresses as base address and stride
+  int stride = 0;
+  ss >> std::hex >> addr;
+  ss >> std::dec >> stride;
+  bool first_bit1_found = false;
+  bool last_bit1_found = false;
+  
+  for (int s = 0; s < WARP_SIZE; s++) {
+    if (!first_bit1_found && mask_bits.test(s)) {
+      first_bit1_found = true;
+      addrs.push_back(addr);
+    } else if (first_bit1_found && !last_bit1_found) {
+      if (mask_bits.test(s)) {
+        addr += stride;
+        addrs.push_back(addr);
+      } else
+        last_bit1_found = true;
+    } else addrs.push_back(0);
+  }
+
+  return addrs;
+}
+
+inline vector<uint64_t> decode_base_delta(
+  std::stringstream &ss,
+  std::bitset<WARP_SIZE> mask_bits
+) {
+  vector<uint64_t> addrs;
+  uint64_t addr = 0;
+  int64_t delta;
+
+  // read addresses as base address and deltas
+  ss >> std::hex >> addr;
+  bool first_bit1_found = false;
+  for(int s = 0; s < WARP_SIZE; s++) {
+    if(!first_bit1_found && mask_bits.test(s)) {
+      first_bit1_found = true;
+      addrs.push_back(addr);
+    } else if(first_bit1_found && mask_bits.test(s)) {
+      ss >> std::dec >> delta;
+      addr += delta;
+      addrs.push_back(addr);
+    } else addrs.push_back(0);
+  }
+
+  return addrs;
+}
\ No newline at end of file
diff --git a/util/tracer_nvbit/tracer_tool/tracer_tool.cu b/util/tracer_nvbit/tracer_tool/tracer_tool.cu
index 5db9ff7eb1..aeca374ee1 100644
--- a/util/tracer_nvbit/tracer_tool/tracer_tool.cu
+++ b/util/tracer_nvbit/tracer_tool/tracer_tool.cu
@@ -39,6 +39,7 @@
 
 /* for general utils functions */
 #include "utils.h"
+#include "decode.h"
 
 #define TRACER_VERSION "5"
 #define TRACER_VERSION_ALLOW_REG_VAL "6"
@@ -252,14 +253,6 @@ bool should_trace_kernel(uint64_t kernel_id, const std::string &kernel_name) {
   return false;
 }
 
-enum address_format {
-  list_all = 0,
-  base_stride = 1,
-  base_delta = 2,
-  tma_list_all = 3,
-  tma_base_delta = 4
-};
-
 /* File pointers for the kernels, and stats files */
 static FILE *kernelsFile = NULL;
 static FILE *statsFile = NULL;
diff --git a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
index 03e1ea4822..aa26182766 100644
--- a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
+++ b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
@@ -23,6 +23,7 @@
 #include <zstd.h>
 
 #include "../tracez_format.h"
+#include "decode.h"
 
 using namespace std;
 
@@ -600,43 +601,9 @@ static mbar_inst_info classify_mbar_inst(const string &line,
   } else {
     unsigned addr_mode;
     ss >> dec >> addr_mode;
-    int first_lane = -1;
-    for (int s = 0; s < 32 && first_lane < 0; s++)
-      if ((mask >> s) & 1u)
-        first_lane = s;
-    if (first_lane < 0)
-      return r;
-    if (addr_mode == 0) { // list_all: one address per active lane
-      for (int s = first_lane; s < 32; s++) {
-        if (!((mask >> s) & 1u))
-          continue;
-        uint64_t addr = 0;
-        ss >> hex >> addr;
-        addrs.push_back(addr);
-      }
-    } else if (addr_mode == 1) { // base_stride: one contiguous run of lanes
-      uint64_t base = 0;
-      long long stride = 0;
-      ss >> hex >> base;
-      ss >> dec >> stride;
-      uint64_t addr = base;
-      for (int s = first_lane; s < 32 && ((mask >> s) & 1u); s++) {
-        addrs.push_back(addr);
-        addr += stride;
-      }
-    } else if (addr_mode == 2) { // base_delta: delta from the previous lane
-      uint64_t addr = 0;
-      ss >> hex >> addr;
-      addrs.push_back(addr);
-      for (int s = first_lane + 1; s < 32; s++) {
-        if (!((mask >> s) & 1u))
-          continue;
-        long long delta = 0;
-        ss >> dec >> delta;
-        addr += delta;
-        addrs.push_back(addr);
-      }
-    } else {
+    try {
+      addrs = decode(addr_mode, ss, mask);
+    } catch (const std::exception&) {
       return r;
     }
   }

From 0735571b8213276427393724753dedd865d3669c Mon Sep 17 00:00:00 2001
From: benlaflen <114254189+benlaflen@users.noreply.github.com>
Date: Tue, 20 Oct 2026 10:42:52 -0400
Subject: [PATCH 2/7] Fix compilation errors related to address decoding
 unification

---
 gpu-simulator/trace-parser/trace_parser.cc    |  2 +-
 gpu-simulator/trace-parser/trace_parser.h     |  5 ++
 util/tracer_nvbit/tracer_tool/decode.h        | 63 +++++++++----------
 .../post-traces-processing.cpp                |  7 ++-
 4 files changed, 41 insertions(+), 36 deletions(-)

diff --git a/gpu-simulator/trace-parser/trace_parser.cc b/gpu-simulator/trace-parser/trace_parser.cc
index 53a7905aa2..c9583cca62 100644
--- a/gpu-simulator/trace-parser/trace_parser.cc
+++ b/gpu-simulator/trace-parser/trace_parser.cc
@@ -243,7 +243,7 @@ bool inst_trace_t::parse_from_string(std::string trace, unsigned trace_version,
     memadd_info->width = get_datawidth_from_opcode(opcode_tokens);
 
     ss >> std::dec >> address_mode;
-    std::vector<uint64_t> addrs = decode(address_mode, ss, mask_bits); //TODO: JUST MAKE SURE THIS GETS IMPORTED
+    std::vector<uint64_t> addrs = decode(address_mode, ss, mask_bits);
     std::memcpy(memadd_info->addrs, addrs.data(), WARP_SIZE * sizeof(uint64_t));
   }
 
diff --git a/gpu-simulator/trace-parser/trace_parser.h b/gpu-simulator/trace-parser/trace_parser.h
index 8627052b1c..9122ee7ce4 100644
--- a/gpu-simulator/trace-parser/trace_parser.h
+++ b/gpu-simulator/trace-parser/trace_parser.h
@@ -32,6 +32,11 @@ enum address_scope {
   SYS_MEM,
 };
 
+enum address_format {
+  tma_list_all = 3,
+  tma_base_delta = 4
+};
+
 struct trace_command {
   std::string command_string;
   command_type m_type;
diff --git a/util/tracer_nvbit/tracer_tool/decode.h b/util/tracer_nvbit/tracer_tool/decode.h
index d22a5331b0..7e95f49108 100644
--- a/util/tracer_nvbit/tracer_tool/decode.h
+++ b/util/tracer_nvbit/tracer_tool/decode.h
@@ -1,41 +1,25 @@
+#pragma once
+
 #include <vector>
 #include <bitset>
+#include <cstdint>
+#include <sstream>
+#include <stdexcept>
 
-#define WARP_SIZE 32
-
-
-enum address_format {
-  list_all = 0,
-  base_stride = 1,
-  base_delta = 2,
-  tma_list_all = 3,
-  tma_base_delta = 4
-};
-
-inline vector<uint64_t> decode(
-    address_format address_mode,
-    std::stringstream &ss,
-    std::bitset<WARP_SIZE> mask_bits
-) {
-  switch(address_mode) {
-    case address_format::list_all: return decode_list_all(ss, mask_bits);
-    case address_format::base_stride: return decode_base_stride(ss, mask_bits);
-    case address_format::base_delta: return decode_base_delta(ss, mask_bits);
-    default: throw std::runtime_error("Unsupported address mode");
-  }
-}
+constexpr int WARP_SIZE = 32
 
-inline vector<uint64_t> decode_list_all(
+inline std::vector<uint64_t> decode_list_all(
   std::stringstream &ss,
-  std::bitset<WARP_SIZE> mask_bits,
+  std::bitset<WARP_SIZE> mask_bits
 ) {
-  vector<uint64_t> addrs;
+  std::vector<uint64_t> addrs;
   uint64_t addr = 0;
 
   // read addresses one by one from the file
   for (int s = 0; s < WARP_SIZE; s++) {
     if (mask_bits.test(s)) {
       ss >> std::hex >> addr;
+      assert(!ss.fail())
       addrs.push_back(addr);
     } else addrs.push_back(0);
   }
@@ -43,17 +27,18 @@ inline vector<uint64_t> decode_list_all(
   return addrs;
 }
 
-inline vector<uint64_t> decode_base_stride(
+inline std::vector<uint64_t> decode_base_stride(
   std::stringstream &ss,
   std::bitset<WARP_SIZE> mask_bits
 ) {
-  vector<uint64_t> addrs;
+  std::vector<uint64_t> addrs;
   uint64_t addr = 0;
 
   // read addresses as base address and stride
   int stride = 0;
   ss >> std::hex >> addr;
   ss >> std::dec >> stride;
+  assert(!ss.fail())
   bool first_bit1_found = false;
   bool last_bit1_found = false;
   
@@ -73,16 +58,17 @@ inline vector<uint64_t> decode_base_stride(
   return addrs;
 }
 
-inline vector<uint64_t> decode_base_delta(
+inline std::vector<uint64_t> decode_base_delta(
   std::stringstream &ss,
   std::bitset<WARP_SIZE> mask_bits
 ) {
-  vector<uint64_t> addrs;
+  std::vector<uint64_t> addrs;
   uint64_t addr = 0;
-  int64_t delta;
+  int64_t delta = 0;
 
   // read addresses as base address and deltas
   ss >> std::hex >> addr;
+  assert(!ss.fail())
   bool first_bit1_found = false;
   for(int s = 0; s < WARP_SIZE; s++) {
     if(!first_bit1_found && mask_bits.test(s)) {
@@ -90,10 +76,23 @@ inline vector<uint64_t> decode_base_delta(
       addrs.push_back(addr);
     } else if(first_bit1_found && mask_bits.test(s)) {
       ss >> std::dec >> delta;
+      assert(!ss.fail())
       addr += delta;
       addrs.push_back(addr);
     } else addrs.push_back(0);
   }
 
   return addrs;
-}
\ No newline at end of file
+}
+
+inline std::vector<uint64_t> decode(
+    unsigned address_mode,
+    std::stringstream &ss,
+    std::bitset<WARP_SIZE> mask_bits
+) {
+
+    if(address_mode == 0) return decode_list_all(ss, mask_bits);
+    if(address_mode == 1) return decode_base_stride(ss, mask_bits);
+    if(address_mode == 2) return decode_base_delta(ss, mask_bits);
+    throw std::runtime_error("Unsupported address mode");
+  }
\ No newline at end of file
diff --git a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
index aa26182766..fe93fc6d83 100644
--- a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
+++ b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
@@ -23,7 +23,7 @@
 #include <zstd.h>
 
 #include "../tracez_format.h"
-#include "decode.h"
+#include "../decode.h"
 
 using namespace std;
 
@@ -601,9 +601,10 @@ static mbar_inst_info classify_mbar_inst(const string &line,
   } else {
     unsigned addr_mode;
     ss >> dec >> addr_mode;
-    try {
+    try {fprintf(stderr, "Warning: %s (mode %u), classifying as OTHER\n", e.what(), addr_mode);
       addrs = decode(addr_mode, ss, mask);
-    } catch (const std::exception&) {
+    } catch (const std::exception &e) {
+      fprintf(stderr, "Warning: %s (mode %u), classifying as OTHER\n", e.what(), addr_mode);
       return r;
     }
   }

From 0d7125abcfceaa0cf82091b2351ed7c7f79f12a0 Mon Sep 17 00:00:00 2001
From: JRPAN <25518778+JRPan@users.noreply.github.com>
Date: Sat, 17 Oct 2026 20:44:38 -0400
Subject: [PATCH 3/7] Run clang-format on the address decoding changes

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
---
 gpu-simulator/trace-parser/trace_parser.h     |  2 +-
 util/tracer_nvbit/tracer_tool/decode.h        | 28 +++++++++++--------
 util/tracer_nvbit/tracer_tool/tracer_tool.cu  |  2 +-
 .../post-traces-processing.cpp                |  2 +-
 4 files changed, 19 insertions(+), 15 deletions(-)

diff --git a/gpu-simulator/trace-parser/trace_parser.h b/gpu-simulator/trace-parser/trace_parser.h
index 9122ee7ce4..060e4da11b 100644
--- a/gpu-simulator/trace-parser/trace_parser.h
+++ b/gpu-simulator/trace-parser/trace_parser.h
@@ -13,7 +13,7 @@
 #include <string>
 #include <vector>
 
-#include "decode.h"
+#include "../../util/tracer_nvbit/tracer_tool/decode.h"
 
 #define MAX_DST 1
 #define MAX_SRC 4
diff --git a/util/tracer_nvbit/tracer_tool/decode.h b/util/tracer_nvbit/tracer_tool/decode.h
index 7e95f49108..0a403ffc38 100644
--- a/util/tracer_nvbit/tracer_tool/decode.h
+++ b/util/tracer_nvbit/tracer_tool/decode.h
@@ -5,8 +5,9 @@
 #include <cstdint>
 #include <sstream>
 #include <stdexcept>
+#include <cassert>
 
-constexpr int WARP_SIZE = 32
+constexpr int WARP_SIZE = 32;
 
 inline std::vector<uint64_t> decode_list_all(
   std::stringstream &ss,
@@ -19,9 +20,10 @@ inline std::vector<uint64_t> decode_list_all(
   for (int s = 0; s < WARP_SIZE; s++) {
     if (mask_bits.test(s)) {
       ss >> std::hex >> addr;
-      assert(!ss.fail())
+      assert(!ss.fail());
       addrs.push_back(addr);
-    } else addrs.push_back(0);
+    } else
+      addrs.push_back(0);
   }
 
   return addrs;
@@ -38,10 +40,10 @@ inline std::vector<uint64_t> decode_base_stride(
   int stride = 0;
   ss >> std::hex >> addr;
   ss >> std::dec >> stride;
-  assert(!ss.fail())
+  assert(!ss.fail());
   bool first_bit1_found = false;
   bool last_bit1_found = false;
-  
+
   for (int s = 0; s < WARP_SIZE; s++) {
     if (!first_bit1_found && mask_bits.test(s)) {
       first_bit1_found = true;
@@ -52,7 +54,8 @@ inline std::vector<uint64_t> decode_base_stride(
         addrs.push_back(addr);
       } else
         last_bit1_found = true;
-    } else addrs.push_back(0);
+    } else
+      addrs.push_back(0);
   }
 
   return addrs;
@@ -68,18 +71,19 @@ inline std::vector<uint64_t> decode_base_delta(
 
   // read addresses as base address and deltas
   ss >> std::hex >> addr;
-  assert(!ss.fail())
+  assert(!ss.fail());
   bool first_bit1_found = false;
-  for(int s = 0; s < WARP_SIZE; s++) {
-    if(!first_bit1_found && mask_bits.test(s)) {
+  for (int s = 0; s < WARP_SIZE; s++) {
+    if (!first_bit1_found && mask_bits.test(s)) {
       first_bit1_found = true;
       addrs.push_back(addr);
-    } else if(first_bit1_found && mask_bits.test(s)) {
+    } else if (first_bit1_found && mask_bits.test(s)) {
       ss >> std::dec >> delta;
-      assert(!ss.fail())
+      assert(!ss.fail());
       addr += delta;
       addrs.push_back(addr);
-    } else addrs.push_back(0);
+    } else
+      addrs.push_back(0);
   }
 
   return addrs;
diff --git a/util/tracer_nvbit/tracer_tool/tracer_tool.cu b/util/tracer_nvbit/tracer_tool/tracer_tool.cu
index aeca374ee1..45f59d3f06 100644
--- a/util/tracer_nvbit/tracer_tool/tracer_tool.cu
+++ b/util/tracer_nvbit/tracer_tool/tracer_tool.cu
@@ -38,8 +38,8 @@
 #include "watchdog.h"
 
 /* for general utils functions */
-#include "utils.h"
 #include "decode.h"
+#include "utils.h"
 
 #define TRACER_VERSION "5"
 #define TRACER_VERSION_ALLOW_REG_VAL "6"
diff --git a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
index fe93fc6d83..be2ce1fc60 100644
--- a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
+++ b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
@@ -601,7 +601,7 @@ static mbar_inst_info classify_mbar_inst(const string &line,
   } else {
     unsigned addr_mode;
     ss >> dec >> addr_mode;
-    try {fprintf(stderr, "Warning: %s (mode %u), classifying as OTHER\n", e.what(), addr_mode);
+    try {
       addrs = decode(addr_mode, ss, mask);
     } catch (const std::exception &e) {
       fprintf(stderr, "Warning: %s (mode %u), classifying as OTHER\n", e.what(), addr_mode);

From d71a20e1bea3c420e20b38b2baca95fd46cdcd90 Mon Sep 17 00:00:00 2001
From: JRPAN <25518778+JRPan@users.noreply.github.com>
Date: Tue, 20 Oct 2026 12:43:56 -0400
Subject: [PATCH 4/7] Apply clang-format

Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
---
 gpu-simulator/trace-parser/trace_parser.h     |  5 +--
 util/tracer_nvbit/tracer_tool/decode.h        | 43 ++++++++-----------
 .../post-traces-processing.cpp                |  5 ++-
 3 files changed, 23 insertions(+), 30 deletions(-)

diff --git a/gpu-simulator/trace-parser/trace_parser.h b/gpu-simulator/trace-parser/trace_parser.h
index 060e4da11b..3f5453fa36 100644
--- a/gpu-simulator/trace-parser/trace_parser.h
+++ b/gpu-simulator/trace-parser/trace_parser.h
@@ -32,10 +32,7 @@ enum address_scope {
   SYS_MEM,
 };
 
-enum address_format {
-  tma_list_all = 3,
-  tma_base_delta = 4
-};
+enum address_format { tma_list_all = 3, tma_base_delta = 4 };
 
 struct trace_command {
   std::string command_string;
diff --git a/util/tracer_nvbit/tracer_tool/decode.h b/util/tracer_nvbit/tracer_tool/decode.h
index 0a403ffc38..c83fae73df 100644
--- a/util/tracer_nvbit/tracer_tool/decode.h
+++ b/util/tracer_nvbit/tracer_tool/decode.h
@@ -1,18 +1,16 @@
 #pragma once
 
-#include <vector>
 #include <bitset>
+#include <cassert>
 #include <cstdint>
 #include <sstream>
 #include <stdexcept>
-#include <cassert>
+#include <vector>
 
 constexpr int WARP_SIZE = 32;
 
-inline std::vector<uint64_t> decode_list_all(
-  std::stringstream &ss,
-  std::bitset<WARP_SIZE> mask_bits
-) {
+inline std::vector<uint64_t> decode_list_all(std::stringstream &ss,
+                                             std::bitset<WARP_SIZE> mask_bits) {
   std::vector<uint64_t> addrs;
   uint64_t addr = 0;
 
@@ -29,10 +27,8 @@ inline std::vector<uint64_t> decode_list_all(
   return addrs;
 }
 
-inline std::vector<uint64_t> decode_base_stride(
-  std::stringstream &ss,
-  std::bitset<WARP_SIZE> mask_bits
-) {
+inline std::vector<uint64_t>
+decode_base_stride(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
   std::vector<uint64_t> addrs;
   uint64_t addr = 0;
 
@@ -61,10 +57,8 @@ inline std::vector<uint64_t> decode_base_stride(
   return addrs;
 }
 
-inline std::vector<uint64_t> decode_base_delta(
-  std::stringstream &ss,
-  std::bitset<WARP_SIZE> mask_bits
-) {
+inline std::vector<uint64_t>
+decode_base_delta(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
   std::vector<uint64_t> addrs;
   uint64_t addr = 0;
   int64_t delta = 0;
@@ -89,14 +83,15 @@ inline std::vector<uint64_t> decode_base_delta(
   return addrs;
 }
 
-inline std::vector<uint64_t> decode(
-    unsigned address_mode,
-    std::stringstream &ss,
-    std::bitset<WARP_SIZE> mask_bits
-) {
+inline std::vector<uint64_t> decode(unsigned address_mode,
+                                    std::stringstream &ss,
+                                    std::bitset<WARP_SIZE> mask_bits) {
 
-    if(address_mode == 0) return decode_list_all(ss, mask_bits);
-    if(address_mode == 1) return decode_base_stride(ss, mask_bits);
-    if(address_mode == 2) return decode_base_delta(ss, mask_bits);
-    throw std::runtime_error("Unsupported address mode");
-  }
\ No newline at end of file
+  if (address_mode == 0)
+    return decode_list_all(ss, mask_bits);
+  if (address_mode == 1)
+    return decode_base_stride(ss, mask_bits);
+  if (address_mode == 2)
+    return decode_base_delta(ss, mask_bits);
+  throw std::runtime_error("Unsupported address mode");
+}
\ No newline at end of file
diff --git a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
index be2ce1fc60..6efa34727f 100644
--- a/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
+++ b/util/tracer_nvbit/tracer_tool/traces-processing/post-traces-processing.cpp
@@ -22,8 +22,8 @@
 #include <zdict.h>
 #include <zstd.h>
 
-#include "../tracez_format.h"
 #include "../decode.h"
+#include "../tracez_format.h"
 
 using namespace std;
 
@@ -604,7 +604,8 @@ static mbar_inst_info classify_mbar_inst(const string &line,
     try {
       addrs = decode(addr_mode, ss, mask);
     } catch (const std::exception &e) {
-      fprintf(stderr, "Warning: %s (mode %u), classifying as OTHER\n", e.what(), addr_mode);
+      fprintf(stderr, "Warning: %s (mode %u), classifying as OTHER\n", e.what(),
+              addr_mode);
       return r;
     }
   }

From 1ced18469bfa3dc84e12efd08af7393ee7ad9d50 Mon Sep 17 00:00:00 2001
From: benlaflen <114254189+benlaflen@users.noreply.github.com>
Date: Mon, 26 Oct 2026 13:43:09 -0400
Subject: [PATCH 5/7] Fixed bug in decode.h where an existing bug in
 trace_parser.cc that read uninitialized memory was faithfully copied for
 accuracy and became breaking

---
 gpu-simulator/trace-parser/trace_parser.h |  2 --
 util/tracer_nvbit/tracer_tool/decode.h    | 25 ++++++++++++++++-------
 2 files changed, 18 insertions(+), 9 deletions(-)

diff --git a/gpu-simulator/trace-parser/trace_parser.h b/gpu-simulator/trace-parser/trace_parser.h
index 3f5453fa36..a6112849ac 100644
--- a/gpu-simulator/trace-parser/trace_parser.h
+++ b/gpu-simulator/trace-parser/trace_parser.h
@@ -32,8 +32,6 @@ enum address_scope {
   SYS_MEM,
 };
 
-enum address_format { tma_list_all = 3, tma_base_delta = 4 };
-
 struct trace_command {
   std::string command_string;
   command_type m_type;
diff --git a/util/tracer_nvbit/tracer_tool/decode.h b/util/tracer_nvbit/tracer_tool/decode.h
index c83fae73df..02c27fdd04 100644
--- a/util/tracer_nvbit/tracer_tool/decode.h
+++ b/util/tracer_nvbit/tracer_tool/decode.h
@@ -9,6 +9,14 @@
 
 constexpr int WARP_SIZE = 32;
 
+enum address_format {
+  list_all = 0,
+  base_stride = 1,
+  base_delta = 2,
+  tma_list_all = 3,
+  tma_base_delta = 4
+};
+
 inline std::vector<uint64_t> decode_list_all(std::stringstream &ss,
                                              std::bitset<WARP_SIZE> mask_bits) {
   std::vector<uint64_t> addrs;
@@ -18,12 +26,12 @@ inline std::vector<uint64_t> decode_list_all(std::stringstream &ss,
   for (int s = 0; s < WARP_SIZE; s++) {
     if (mask_bits.test(s)) {
       ss >> std::hex >> addr;
-      assert(!ss.fail());
+      if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
       addrs.push_back(addr);
     } else
       addrs.push_back(0);
   }
-
+  assert(addrs.size() == WARP_SIZE);
   return addrs;
 }
 
@@ -36,7 +44,7 @@ decode_base_stride(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
   int stride = 0;
   ss >> std::hex >> addr;
   ss >> std::dec >> stride;
-  assert(!ss.fail());
+  if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
   bool first_bit1_found = false;
   bool last_bit1_found = false;
 
@@ -50,10 +58,11 @@ decode_base_stride(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
         addrs.push_back(addr);
       } else
         last_bit1_found = true;
+        addrs.push_back(0);
     } else
       addrs.push_back(0);
   }
-
+  assert(addrs.size() == WARP_SIZE);
   return addrs;
 }
 
@@ -65,7 +74,7 @@ decode_base_delta(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
 
   // read addresses as base address and deltas
   ss >> std::hex >> addr;
-  assert(!ss.fail());
+  if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
   bool first_bit1_found = false;
   for (int s = 0; s < WARP_SIZE; s++) {
     if (!first_bit1_found && mask_bits.test(s)) {
@@ -73,13 +82,13 @@ decode_base_delta(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
       addrs.push_back(addr);
     } else if (first_bit1_found && mask_bits.test(s)) {
       ss >> std::dec >> delta;
-      assert(!ss.fail());
+      if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
       addr += delta;
       addrs.push_back(addr);
     } else
       addrs.push_back(0);
   }
-
+  assert(addrs.size() == WARP_SIZE);
   return addrs;
 }
 
@@ -93,5 +102,7 @@ inline std::vector<uint64_t> decode(unsigned address_mode,
     return decode_base_stride(ss, mask_bits);
   if (address_mode == 2)
     return decode_base_delta(ss, mask_bits);
+  if(address_mode == 3 || address_mode == 4)
+    throw std::runtime_error("decode() does not support TMA addressing modes as they are written usage-specific; reference usage-specific code at other callsites (for instance, trace_parser.cc or post-traces-processing.cpp) and do not invoke decode() here");
   throw std::runtime_error("Unsupported address mode");
 }
\ No newline at end of file

From 06020320f8645d98be14aabf1f0194e58995eff3 Mon Sep 17 00:00:00 2001
From: Benjamin Laflen <blaflen@tgrogers-gpu01.ecn.purdue.edu>
Date: Tue, 27 Oct 2026 12:15:58 -0400
Subject: [PATCH 6/7] Ran formatter tool

---
 util/tracer_nvbit/tracer_tool/decode.h | 22 +++++++++++++++-------
 1 file changed, 15 insertions(+), 7 deletions(-)

diff --git a/util/tracer_nvbit/tracer_tool/decode.h b/util/tracer_nvbit/tracer_tool/decode.h
index 02c27fdd04..ea00b99d4a 100644
--- a/util/tracer_nvbit/tracer_tool/decode.h
+++ b/util/tracer_nvbit/tracer_tool/decode.h
@@ -26,7 +26,8 @@ inline std::vector<uint64_t> decode_list_all(std::stringstream &ss,
   for (int s = 0; s < WARP_SIZE; s++) {
     if (mask_bits.test(s)) {
       ss >> std::hex >> addr;
-      if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
+      if (ss.fail())
+        throw std::runtime_error("String stream input failed prematurely");
       addrs.push_back(addr);
     } else
       addrs.push_back(0);
@@ -44,7 +45,8 @@ decode_base_stride(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
   int stride = 0;
   ss >> std::hex >> addr;
   ss >> std::dec >> stride;
-  if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
+  if (ss.fail())
+    throw std::runtime_error("String stream input failed prematurely");
   bool first_bit1_found = false;
   bool last_bit1_found = false;
 
@@ -58,7 +60,7 @@ decode_base_stride(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
         addrs.push_back(addr);
       } else
         last_bit1_found = true;
-        addrs.push_back(0);
+      addrs.push_back(0);
     } else
       addrs.push_back(0);
   }
@@ -74,7 +76,8 @@ decode_base_delta(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
 
   // read addresses as base address and deltas
   ss >> std::hex >> addr;
-  if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
+  if (ss.fail())
+    throw std::runtime_error("String stream input failed prematurely");
   bool first_bit1_found = false;
   for (int s = 0; s < WARP_SIZE; s++) {
     if (!first_bit1_found && mask_bits.test(s)) {
@@ -82,7 +85,8 @@ decode_base_delta(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
       addrs.push_back(addr);
     } else if (first_bit1_found && mask_bits.test(s)) {
       ss >> std::dec >> delta;
-      if(ss.fail()) throw std::runtime_error("String stream input failed prematurely");
+      if (ss.fail())
+        throw std::runtime_error("String stream input failed prematurely");
       addr += delta;
       addrs.push_back(addr);
     } else
@@ -102,7 +106,11 @@ inline std::vector<uint64_t> decode(unsigned address_mode,
     return decode_base_stride(ss, mask_bits);
   if (address_mode == 2)
     return decode_base_delta(ss, mask_bits);
-  if(address_mode == 3 || address_mode == 4)
-    throw std::runtime_error("decode() does not support TMA addressing modes as they are written usage-specific; reference usage-specific code at other callsites (for instance, trace_parser.cc or post-traces-processing.cpp) and do not invoke decode() here");
+  if (address_mode == 3 || address_mode == 4)
+    throw std::runtime_error(
+        "decode() does not support TMA addressing modes as they are written "
+        "usage-specific; reference usage-specific code at other callsites (for "
+        "instance, trace_parser.cc or post-traces-processing.cpp) and do not "
+        "invoke decode() here");
   throw std::runtime_error("Unsupported address mode");
 }
\ No newline at end of file

From f8f8cea7b544e79e3c974e052f5d148c40e7d6ac Mon Sep 17 00:00:00 2001
From: benlaflen <114254189+benlaflen@users.noreply.github.com>
Date: Tue, 27 Oct 2026 12:38:16 -0400
Subject: [PATCH 7/7] Fix a bug related to a missing close brace

---
 util/tracer_nvbit/tracer_tool/decode.h | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/util/tracer_nvbit/tracer_tool/decode.h b/util/tracer_nvbit/tracer_tool/decode.h
index ea00b99d4a..b475281761 100644
--- a/util/tracer_nvbit/tracer_tool/decode.h
+++ b/util/tracer_nvbit/tracer_tool/decode.h
@@ -58,9 +58,10 @@ decode_base_stride(std::stringstream &ss, std::bitset<WARP_SIZE> mask_bits) {
       if (mask_bits.test(s)) {
         addr += stride;
         addrs.push_back(addr);
-      } else
+      } else {
         last_bit1_found = true;
-      addrs.push_back(0);
+        addrs.push_back(0);
+      }
     } else
       addrs.push_back(0);
   }