Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
atgutier merged 3 commits into from
Aug 19, 2026
Merged

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

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
09908d1
[rocjitsu] Refactor test encoding handling in execution tests (NFC)
amd-eochoalo Aug 19, 2026
b2016bb
Add missing header
amd-eochoalo Aug 19, 2026
2c5f682
Use pragma and further refactoring
amd-eochoalo Aug 19, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
39 changes: 12 additions & 27 deletions emulation/rocjitsu/tests/decode_execute_benchmark.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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 {
Expand Down Expand Up @@ -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 test_encodings::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;

Expand Down Expand Up @@ -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");
Expand Down Expand Up @@ -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
20 changes: 5 additions & 15 deletions emulation/rocjitsu/tests/instruction_execution_harness_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down Expand Up @@ -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);
}
Expand Down Expand Up @@ -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 test_encodings::mnemonic_has_any_prefix(mn, SKIP_PREFIXES);
}

constexpr std::array<std::string_view, 1> EXPECTED_CDNA_UNIMPLEMENTED = {"s_setvskip"};
Expand Down Expand Up @@ -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");
Expand Down Expand Up @@ -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"); }
Expand Down
27 changes: 27 additions & 0 deletions emulation/rocjitsu/tests/test_encodings_util.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,27 @@
// Copyright (c) 2026 Advanced Micro Devices, Inc.
// SPDX-License-Identifier: MIT

#ifndef ROCJITSU_TESTS_TEST_ENCODINGS_UTIL_H_
#define ROCJITSU_TESTS_TEST_ENCODINGS_UTIL_H_
Comment thread
amd-eochoalo marked this conversation as resolved.
Resolved

#include <cstddef>
#include <string_view>

namespace rocjitsu::test_encodings {

inline bool mnemonic_has_prefix(std::string_view mnemonic, std::string_view prefix) {
return mnemonic.size() >= prefix.size() && mnemonic.substr(0, prefix.size()) == prefix;
}

template <size_t N>
inline bool mnemonic_has_any_prefix(std::string_view mnemonic,
const std::string_view (&prefixes)[N]) {
for (auto prefix : prefixes)
if (mnemonic_has_prefix(mnemonic, prefix))
return true;
return false;
}

} // namespace rocjitsu::test_encodings

#endif // ROCJITSU_TESTS_TEST_ENCODINGS_UTIL_H_
Loading