[rocjitsu] Refactor test encoding handling in execution tests (NFC)

emulation/rocjitsu/tests/decode_execute_benchmark.cpp

@@ -19,6 +19,7 @@
 #include "rocjitsu/vm/amdgpu/gpu_memory.h"
 #include "rocjitsu/vm/amdgpu/l2_cache.h"
 #include "rocjitsu/vm/amdgpu/wavefront.h"
+#include "test_encodings_util.h"
 #include "util/except.h"
 
 #include "rocjitsu/isa/arch/amdgpu/generated/cdna4/test_encodings.h"
@@ -39,12 +40,6 @@ namespace {
 using namespace rocjitsu;
 using Clock = std::chrono::steady_clock;
 
-/// @brief Common test encoding entry (matches all ISA test_encodings.h).
-struct TestEncEntry {
-  std::string_view mnemonic;
-  std::array<uint32_t, 2> words;
-};
-
 enum class InstCategory : uint8_t { SCALAR_ALU, VECTOR_ALU, MEMORY };
 
 struct CorpusEntry {
@@ -75,15 +70,13 @@ inline bool should_skip(std::string_view mn) {
       "s_icache_inv", "s_sendmsghalt", "v_readlane", "v_writelane", "v_readfirstlane", "v_mfma_",
       "v_permlane",   "v_smfma_",      "v_wmma_",    "v_swmmac_",
   };
-  for (auto p : SKIP)
-    if (mn.substr(0, p.size()) == p)
-      return true;
-  return false;
+  return mnemonic_has_any_prefix(mn, SKIP);
 }
 
 /// @brief Build a corpus of decodable instructions from test data.
 /// Only filters by skip list and decodability — does not try execute.
-std::vector<CorpusEntry> build_corpus(const TestEncEntry *encodings, size_t num_encodings,
+template <typename Encoding>
+std::vector<CorpusEntry> build_corpus(const Encoding *encodings, size_t num_encodings,
                                       Decoder &decoder) {
   std::vector<CorpusEntry> corpus;
 
@@ -116,7 +109,8 @@ struct TimingResult {
 };
 
 /// @brief Run the benchmark for a given ISA.
-void run_benchmark(rj_code_arch_t arch, std::string_view arch_name, const TestEncEntry *encodings,
+template <typename Encoding>
+void run_benchmark(rj_code_arch_t arch, std::string_view arch_name, const Encoding *encodings,
                    size_t num_encodings) {
   // Set up CU + wavefront.
   amdgpu::GpuMemory gpu_mem(std::string(arch_name) + "_bench_mem");
@@ -242,22 +236,13 @@ void run_benchmark(rj_code_arch_t arch, std::string_view arch_name, const TestEn
 
 // --- Benchmark tests ---
 
-TEST(DecodeExecuteBenchmark, Cdna4) {
-  run_benchmark(ROCJITSU_CODE_ARCH_CDNA4, "cdna4",
-                reinterpret_cast<const TestEncEntry *>(cdna4::test_data::ENCODINGS),
-                cdna4::test_data::NUM_ENCODINGS);
-}
+#define RUN_BENCHMARK(ISA_NS, ARCH_ENUM, ARCH_STR)                                                 \
+  run_benchmark(ARCH_ENUM, ARCH_STR, ISA_NS::test_data::ENCODINGS, ISA_NS::test_data::NUM_ENCODINGS)
 
-TEST(DecodeExecuteBenchmark, Rdna4) {
-  run_benchmark(ROCJITSU_CODE_ARCH_RDNA4, "rdna4",
-                reinterpret_cast<const TestEncEntry *>(rdna4::test_data::ENCODINGS),
-                rdna4::test_data::NUM_ENCODINGS);
-}
+TEST(DecodeExecuteBenchmark, Cdna4) { RUN_BENCHMARK(cdna4, ROCJITSU_CODE_ARCH_CDNA4, "cdna4"); }
 
-TEST(DecodeExecuteBenchmark, Gfx1250) {
-  run_benchmark(ROCJITSU_CODE_ARCH_CDNA5, "gfx1250",
-                reinterpret_cast<const TestEncEntry *>(cdna5::test_data::ENCODINGS),
-                cdna5::test_data::NUM_ENCODINGS);
-}
+TEST(DecodeExecuteBenchmark, Rdna4) { RUN_BENCHMARK(rdna4, ROCJITSU_CODE_ARCH_RDNA4, "rdna4"); }
+
+TEST(DecodeExecuteBenchmark, Gfx1250) { RUN_BENCHMARK(cdna5, ROCJITSU_CODE_ARCH_CDNA5, "gfx1250"); }
 
 } // namespace

emulation/rocjitsu/tests/instruction_execution_harness_test.cpp

@@ -34,6 +34,7 @@
 #include "rocjitsu/vm/amdgpu/hwreg.h"
 #include "rocjitsu/vm/amdgpu/l2_cache.h"
 #include "rocjitsu/vm/amdgpu/wavefront.h"
+#include "test_encodings_util.h"
 #include "util/data_types.h"
 #include "util/except.h"
 #include "util/simd.h"
@@ -74,12 +75,6 @@ namespace {
 
 using namespace rocjitsu;
 
-/// @brief Common test encoding entry (matches all ISA test_encodings.h).
-struct TestEncEntry {
-  std::string_view mnemonic;
-  std::array<uint32_t, 2> words;
-};
-
 constexpr uint32_t pack16(uint16_t lo, uint16_t hi) {
   return static_cast<uint32_t>(lo) | (static_cast<uint32_t>(hi) << 16);
 }
@@ -237,10 +232,7 @@ inline bool should_skip_inst(std::string_view mn) {
       "s_setreg",
       "s_rfe",
   };
-  for (auto p : SKIP_PREFIXES)
-    if (mn.substr(0, p.size()) == p)
-      return true;
-  return false;
+  return mnemonic_has_any_prefix(mn, SKIP_PREFIXES);
 }
 
 constexpr std::array<std::string_view, 1> EXPECTED_CDNA_UNIMPLEMENTED = {"s_setvskip"};
@@ -271,8 +263,9 @@ const HarnessExpectation *harness_expectation(std::string_view arch_name) {
 }
 
 /// @brief Helper to run all test encodings for a given ISA.
+template <typename Encoding>
 void run_execution_harness(rj_code_arch_t arch, std::string_view arch_name,
-                           const TestEncEntry *encodings, size_t num_encodings) {
+                           const Encoding *encodings, size_t num_encodings) {
   // Set up minimal CU with zeroed memory.
   amdgpu::GpuMemory gpu_mem(std::string(arch_name) + "_mem");
   amdgpu::L2Cache l2(std::string(arch_name) + "_l2");
@@ -387,11 +380,8 @@ void run_execution_harness(rj_code_arch_t arch, std::string_view arch_name,
 
 // --- Parameterized tests per ISA ---
 
-// Helper macro: cast the ISA-specific TestEncoding array to the common layout.
-// The struct layout (string_view + array<uint32_t,2>) is identical across all ISAs.
 #define RUN_HARNESS(ISA_NS, ARCH_ENUM, ARCH_STR)                                                   \
-  run_execution_harness(ARCH_ENUM, ARCH_STR,                                                       \
-                        reinterpret_cast<const TestEncEntry *>(ISA_NS::test_data::ENCODINGS),      \
+  run_execution_harness(ARCH_ENUM, ARCH_STR, ISA_NS::test_data::ENCODINGS,                         \
                         ISA_NS::test_data::NUM_ENCODINGS)
 
 TEST(InstructionExecutionHarness, Cdna1) { RUN_HARNESS(cdna1, ROCJITSU_CODE_ARCH_CDNA1, "cdna1"); }

emulation/rocjitsu/tests/test_encodings_util.h

@@ -0,0 +1,20 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <cstddef>
+#include <string_view>
+
+namespace rocjitsu {
+
+template <std::size_t EncodingCount>
+inline bool mnemonic_has_any_prefix(std::string_view mnemonic,
+                                    const std::string_view (&prefixes)[EncodingCount]) {
+  for (std::string_view prefix : prefixes)
+    if (mnemonic.starts_with(prefix))
+      return true;
+  return false;
+}
+
+} // namespace rocjitsu