diff --git a/emulation/rocjitsu/docs/race-detector.md b/emulation/rocjitsu/docs/race-detector.md index 211d0c2798a8..29906cdd324c 100644 --- a/emulation/rocjitsu/docs/race-detector.md +++ b/emulation/rocjitsu/docs/race-detector.md @@ -148,8 +148,10 @@ races. Some examples: - **VGPR races**: a vector register is read or overwritten by an instruction before a pending global or LDS load has completed (`s_waitcnt vmcnt` / `s_waitcnt lgkmcnt` insufficient). -- **SGPR races**: a scalar register is read before a pending scalar load has - completed (`s_waitcnt lgkmcnt` insufficient). +- **Scalar-register races**: an SGPR or TTMP is read or overwritten before a + pending scalar-memory load has completed. A later scalar load to the same + destination is also checked because scalar-memory results can complete out + of order. - **LDS races**: an LDS byte is read by one wave while another wave has an outstanding write to the same byte, or written by one wave while another wave has an outstanding read of the same byte, without an intervening @@ -183,9 +185,15 @@ operations are in flight. When an instruction in the emulator accesses an LDS byte, there is a check to see what memory events are still in flight that read/write that byte, from the perspective of the accessing thread. In this way, RAW (read-after-write) and WAR (write-after-read) hazards can be detected. -For VGPRs, the detector also flags WAW when an instruction write can be -clobbered by a pending asynchronous load. Similar logic applies for VGPR and -SGPR reads. +For VGPRs and scalar registers, the detector also flags WAW when an instruction +write can be clobbered by a pending asynchronous load. The detector also reports +WAW between scalar loads targeting the same SGPR or TTMP. + +On architectures where scalar-memory and data-share operations use a combined +`lgkmcnt`, a nonzero partial wait cannot identify which scalar destination has +completed. Those scalar destinations remain pending until `lgkmcnt(0)`. The +same partial wait can still retire older data-share operations that are provably +complete from their in-order completion rule. **LDS race detection** uses coarse-grained counters (one per 16-byte chunk) for fast-path checks, with interval-based overlap scanning as a fallback. Live @@ -298,8 +306,10 @@ ctest --test-dir $BUILD_DIR -R "RaceTest" races, races between dispatches, or host-device synchronization issues. - **Limited WAW detection**: VGPR WAW is limited to synchronous instruction - writes that overlap a pending asynchronous load. WAW between two asynchronous - VGPR writers and LDS WAW are not currently reported. + writes that overlap a pending asynchronous load. Scalar-register WAW covers + instruction writes and scalar loads that overlap pending scalar-memory + destinations. WAW between two asynchronous VGPR writers and LDS WAW are not + currently reported. - **Conservative DPP/SDWA write masks**: WAW precision depends on the execution plugin's instruction-write lane and byte masks. DPP destinations can currently diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp index 6de8d3393c6e..96e77c179c84 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp @@ -60,6 +60,8 @@ void WaveRaceState::registerScalarLoad(uint64_t pc, RegisterRef destination, uin if (destination.index >= limit || destination.width > limit - destination.index) return; + checkScalarWrite(destination); + std::vector registers(destination.width); for (uint32_t i = 0; i < destination.width; ++i) registers[i] = destination.index + i; @@ -297,7 +299,7 @@ bool WaveRaceState::isOutstandingFromVgpr(int lane, int reg) const { return false; } -void WaveRaceState::checkScalarRead(RegisterRef ref) const { +void WaveRaceState::checkScalarAccess(RegisterRef ref, bool isWrite) const { const std::vector> *events = nullptr; const std::vector *counts = nullptr; RaceViolation::Space space; @@ -318,16 +320,28 @@ void WaveRaceState::checkScalarRead(RegisterRef ref) const { if (ref.width == 0 || ref.index >= events->size() || ref.width > events->size() - ref.index) return; + std::vector reportedEvents; for (uint32_t offset = 0; offset < ref.width; ++offset) { const uint32_t index = ref.index + offset; if ((*counts)[index] == 0) continue; for (EventId eid : (*events)[index]) { - if (detector->events().type(eid) == type) - detector->getRaceHandler()({space, static_cast(index), waveId.value, -1, false, - detector->getWorkgroupId(), eid}); + if (detector->events().type(eid) != type || + std::find(reportedEvents.begin(), reportedEvents.end(), eid) != reportedEvents.end()) + continue; + reportedEvents.push_back(eid); + detector->getRaceHandler()({space, static_cast(index), waveId.value, -1, isWrite, + detector->getWorkgroupId(), eid}); } } } +void WaveRaceState::checkScalarRead(RegisterRef ref) const { + checkScalarAccess(ref, /*isWrite=*/false); +} + +void WaveRaceState::checkScalarWrite(RegisterRef ref) const { + checkScalarAccess(ref, /*isWrite=*/true); +} + } // namespace rocjitsu::plugins::race_detector diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h index 3745b4a3cb75..97f8e99538b8 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h @@ -94,6 +94,10 @@ class WaveRaceState { /// violation (outstanding scalar load targeting the same register). void checkScalarRead(RegisterRef reg) const; + /// Check a scalar instruction write for conflicts with pending scalar loads. + /// A wide write reports each conflicting memory event at most once. + void checkScalarWrite(RegisterRef reg) const; + /// True if any outstanding global/LDS store reads from the given VGPR lane. bool isOutstandingFromVgpr(int lane, int reg) const; @@ -120,6 +124,7 @@ class WaveRaceState { uint64_t execMask, uint8_t byteMask, IntervalSet ldsIntervals, amdgpu::WaitCounterType waitCounterType); void retireEventRegisters(EventId); + void checkScalarAccess(RegisterRef reg, bool isWrite) const; template void resolveWaitCnt(int limit, Pred isTargetType); void applyWaitCounter(amdgpu::WaitCounterType type, int threshold); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp index afa7e25a137e..331c79dac713 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp @@ -378,6 +378,12 @@ void RaceDetectorPlugin::onAmdgpuReadScalarRegister(const amdgpu::Wavefront *wf, s->race_state->checkScalarRead(reg); } +void RaceDetectorPlugin::onAmdgpuWriteScalarRegister(const amdgpu::Wavefront *wf, RegisterRef reg) { + auto *s = get_state(wf); + assert(s && s->race_state); + s->race_state->checkScalarWrite(reg); +} + void RaceDetectorPlugin::onAmdgpuBeforeExecuteInstruction(uint64_t pc, const Instruction &inst, amdgpu::Wavefront &wf) { auto *s = get_state(wf); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h index a92312dee27a..889293183b0b 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h @@ -113,6 +113,8 @@ class RaceDetectorPlugin : public ExecutionPlugin { void onAmdgpuReadScalarRegister(const amdgpu::Wavefront *wf, RegisterRef reg) override; + void onAmdgpuWriteScalarRegister(const amdgpu::Wavefront *wf, RegisterRef reg) override; + void onAmdgpuBeforeExecuteInstruction(uint64_t pc, const Instruction &inst, amdgpu::Wavefront &wf) override; diff --git a/emulation/rocjitsu/tests/execution_plugin_test.cpp b/emulation/rocjitsu/tests/execution_plugin_test.cpp index aa01a9f16067..c1bb6df2e511 100644 --- a/emulation/rocjitsu/tests/execution_plugin_test.cpp +++ b/emulation/rocjitsu/tests/execution_plugin_test.cpp @@ -2951,6 +2951,35 @@ TEST(RaceDetectorPluginTest, ScalarLoadToTtmpReportsReadBeforeWait) { EXPECT_NE(sink.str().find("type=TTMP"), std::string::npos); } +TEST(RaceDetectorPluginTest, ScalarLoadToTtmpReportsWriteBeforeWait) { + PluginFixture f(/*num_wf_slots=*/1); + PluginSinkConfig sink_config; + StringSink &sink = sink_config.emplace(); + f.plugin_group_ = std::make_shared(std::move(sink_config)); + ASSERT_TRUE(f.plugin_group_->add(std::make_unique())); + f.soc->set_plugin_group(f.plugin_group_); + f.plugin_group_->onInit(); + + auto *cu = f.cu(); + auto *wf = cu->dispatch_wf(/*wg_id=*/0, /*pc=*/0, /*sgprs=*/104, /*vgprs=*/256); + ASSERT_NE(wf, nullptr); + wf->set_exec(1u); + std::array waves{wf}; + f.plugin_group_->onAmdgpuWorkgroupDispatched( + /*dispatch_id=*/1, /*wg_id=*/0, /*physical_vgpr_count=*/256, + /*physical_sgpr_count=*/104, waves); + + auto state = std::make_unique(); + state->dst_register = {ScalarRegisterStorage::TTMP, 0, 1}; + state->is_load = true; + TestMemoryInstruction load(std::move(state)); + f.plugin_group_->onAmdgpuRouteMemoryInstruction(load, *wf); + + RegisterAccess(*wf).write_ttmp(0, 0x12345678u); + EXPECT_NE(sink.str().find("type=TTMP"), std::string::npos); + EXPECT_NE(sink.str().find("access=write"), std::string::npos); +} + TEST(RaceDetectorPluginTest, ScalarLoadToTtmpHonorsSplitKmcntWait) { PluginFixture f(/*num_wf_slots=*/1, /*arch=*/"rdna4", /*wavefront_size=*/32, /*sgprs_per_wf=*/128); diff --git a/emulation/rocjitsu/tests/race-detector/CMakeLists.txt b/emulation/rocjitsu/tests/race-detector/CMakeLists.txt index d10d15ebeac9..67b66f3d9a30 100644 --- a/emulation/rocjitsu/tests/race-detector/CMakeLists.txt +++ b/emulation/rocjitsu/tests/race-detector/CMakeLists.txt @@ -15,13 +15,23 @@ set(RJ_RACE_TEST_SUPPORT_HEADERS ${CMAKE_CURRENT_SOURCE_DIR}/race_log_expectation.hpp ${CMAKE_CURRENT_SOURCE_DIR}/race_test_support.hpp ) -rj_add_hip_test( + +function(rj_add_hip_race_test TEST_NAME) + rj_add_hip_test(${TEST_NAME} ${ARGN}) + # The race plugin is loaded at runtime, so no link dependency pulls its + # shared object into focused builds of the HIP race-test targets. + if(TARGET rocjitsu_plugin_race_so) + add_dependencies(${TEST_NAME}_target rocjitsu_plugin_race_so) + endif() +endfunction() + +rj_add_hip_race_test( hip_race_tests_gfx950 ARCH gfx950 SOURCE hip_race_gfx950_test.hip DEPENDENCIES ${RJ_RACE_TEST_SUPPORT_HEADERS} ) -rj_add_hip_test( +rj_add_hip_race_test( hip_race_tests_gfx1151 ARCH gfx1151 SOURCE hip_race_gfx1151_test.hip @@ -132,6 +142,10 @@ rj_add_gfx950_race_test_case(vgpr_waitcnt) rj_add_gfx950_race_test_case(vgpr_waitcnt_race) rj_add_gfx950_race_test_case(sgpr_waitcnt) rj_add_gfx950_race_test_case(sgpr_waitcnt_race) +rj_add_gfx950_race_test_case(sgpr_waw_load_then_mov) +rj_add_gfx950_race_test_case(sgpr_waw_load_then_mov_race) +rj_add_gfx950_race_test_case(sgpr_waw_load_then_load) +rj_add_gfx950_race_test_case(sgpr_waw_load_then_load_race) rj_add_gfx950_race_test_case(lds_cross_wave) rj_add_gfx950_race_test_case(lds_cross_wave_race) rj_add_gfx950_race_test_case(lds_same_wave_order) diff --git a/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip b/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip index 162ae8cf0234..30501aa9725e 100644 --- a/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip +++ b/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip @@ -128,6 +128,103 @@ TEST_F(Gfx950RaceTest, sgpr_waitcnt_race) { ExpectRace(kSgprWaitcntRaceExpectation); } +// --- SGPR write-after-write --- + +__global__ void sgpr_waw_load_then_mov_kernel(const int *src, int *dst) { + int val; + asm volatile( + "s_load_dword %0, %1, 0x0\n" + "s_waitcnt lgkmcnt(0)\n" + "s_mov_b32 %0, 7\n" + : "=&s"(val) : "s"(src) : "memory" + ); + dst[threadIdx.x] = val; +} + +static constexpr RaceExpectation kSgprWawLoadThenMovExpectation{ + .kernel = "sgpr_waw_load_then_mov_race_kernel", + .type = "SGPR", + .access = "write", + .wave = 0, + .producer = "s_load_dword", + .consumer = "s_mov_b32", +}; + +__global__ void sgpr_waw_load_then_mov_race_kernel(const int *src, int *dst) { + int val; + asm volatile( + "s_load_dword %0, %1, 0x0\n" + // BUG: no wait before a scalar instruction overwrites the load destination. + "s_mov_b32 %0, 7\n" + : "=&s"(val) : "s"(src) + ); + asm volatile("s_waitcnt lgkmcnt(0)\n" ::: "memory"); + dst[threadIdx.x] = val; +} + +static constexpr RaceExpectation kSgprWawLoadThenLoadExpectation{ + .kernel = "sgpr_waw_load_then_load_race_kernel", + .type = "SGPR", + .access = "write", + .wave = 0, + .producer = "s_load_dword", + .consumer = "s_load_dword", +}; + +__global__ void sgpr_waw_load_then_load_kernel(const int *src0, const int *src1, int *dst) { + int val; + asm volatile( + "s_load_dword %0, %1, 0x0\n" + "s_waitcnt lgkmcnt(0)\n" + "s_load_dword %0, %2, 0x0\n" + : "=&s"(val) : "s"(src0), "s"(src1) : "memory" + ); + asm volatile("s_waitcnt lgkmcnt(0)\n" ::: "memory"); + dst[threadIdx.x] = val; +} + +__global__ void sgpr_waw_load_then_load_race_kernel(const int *src0, const int *src1, int *dst) { + int val; + asm volatile( + "s_load_dword %0, %1, 0x0\n" + // BUG: scalar-memory reads can return out of order. + "s_load_dword %0, %2, 0x0\n" + : "=&s"(val) : "s"(src0), "s"(src1) + ); + asm volatile("s_waitcnt lgkmcnt(0)\n" ::: "memory"); + dst[threadIdx.x] = val; +} + +TEST_F(Gfx950RaceTest, sgpr_waw_load_then_mov) { + auto *src = allocWithData(64), *dst = alloc(64); + sgpr_waw_load_then_mov_kernel<<<1, 64>>>(src, dst); + sync(); + ExpectNoRace(); +} + +TEST_F(Gfx950RaceTest, sgpr_waw_load_then_mov_race) { + auto *src = allocWithData(64), *dst = alloc(64); + sgpr_waw_load_then_mov_race_kernel<<<1, 64>>>(src, dst); + sync(); + ExpectRace(kSgprWawLoadThenMovExpectation); +} + +TEST_F(Gfx950RaceTest, sgpr_waw_load_then_load) { + auto *src0 = allocWithData(64), *src1 = allocWithData(64); + auto *dst = alloc(64); + sgpr_waw_load_then_load_kernel<<<1, 64>>>(src0, src1, dst); + sync(); + ExpectNoRace(); +} + +TEST_F(Gfx950RaceTest, sgpr_waw_load_then_load_race) { + auto *src0 = allocWithData(64), *src1 = allocWithData(64); + auto *dst = alloc(64); + sgpr_waw_load_then_load_race_kernel<<<1, 64>>>(src0, src1, dst); + sync(); + ExpectRace(kSgprWawLoadThenLoadExpectation); +} + // --- LDS cross-wave --- __attribute__((amdgpu_flat_work_group_size(128, 128))) diff --git a/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp b/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp index abc7a7f85fa9..453674cf069d 100644 --- a/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp +++ b/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp @@ -9,6 +9,8 @@ // Test categories: // Vgpr_* — VGPR races from global loads (vmcnt) // Sgpr_* — SGPR races from scalar loads (lgkmcnt) +// SgprWaw_* — writes overlapping pending SGPR destinations +// TtmpWaw_* — writes overlapping pending TTMP destinations // LdsCrossWave_* — cross-wave LDS races (missing barrier) // LdsSameWave_* — same-wave LDS instruction ordering // SameWave_* — same-wave VGPR races via LDS loads @@ -226,6 +228,141 @@ TEST(RaceDetector, RejectsOutOfRangeScalarLoadDestination) { EXPECT_EQ(b.events().totalAllocated(), 0); } +// ---- Scalar write-after-write races ---- + +TEST(RaceDetector, SgprWaw_ScalarLoadThenInstructionWrite) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/4, /*numRegs=*/1); + b.checkSgprWrite(/*wave=*/0, /*reg=*/4); + + ASSERT_EQ(b.raceCount(), 1); + const auto &violation = b.violations().front(); + EXPECT_EQ(violation.space, RaceViolation::Space::SGPR); + EXPECT_EQ(violation.index, 4); + EXPECT_EQ(violation.lane, -1); + EXPECT_TRUE(violation.isWrite); + EXPECT_EQ(violation.conflictingEvent, EventId{0}); +} + +TEST(RaceDetector, TtmpWaw_ScalarLoadThenInstructionWrite) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.ttmpLoad(/*wave=*/0, /*ttmpBase=*/4, /*numRegs=*/1); + b.checkTtmpWrite(/*wave=*/0, /*reg=*/4); + + ASSERT_EQ(b.raceCount(), 1); + const auto &violation = b.violations().front(); + EXPECT_EQ(violation.space, RaceViolation::Space::TTMP); + EXPECT_EQ(violation.index, 4); + EXPECT_EQ(violation.lane, -1); + EXPECT_TRUE(violation.isWrite); + EXPECT_EQ(violation.conflictingEvent, EventId{0}); +} + +TEST(RaceDetector, SgprWaw_ScalarLoadThenScalarLoad) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/4, /*numRegs=*/1); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/4, /*numRegs=*/1); + + ASSERT_EQ(b.raceCount(), 1); + EXPECT_TRUE(b.hasSgprRace(4)); + EXPECT_TRUE(b.violations().front().isWrite); + EXPECT_EQ(b.violations().front().conflictingEvent, EventId{0}); +} + +TEST(RaceDetector, TtmpWaw_ScalarLoadThenScalarLoad) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.ttmpLoad(/*wave=*/0, /*ttmpBase=*/4, /*numRegs=*/1); + b.ttmpLoad(/*wave=*/0, /*ttmpBase=*/4, /*numRegs=*/1); + + ASSERT_EQ(b.raceCount(), 1); + EXPECT_TRUE(b.hasTtmpRace(4)); + EXPECT_TRUE(b.violations().front().isWrite); + EXPECT_EQ(b.violations().front().conflictingEvent, EventId{0}); +} + +TEST(RaceDetector, SgprWaw_WaitClearsPendingLoad) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/4, /*numRegs=*/1); + b.waitcnt(/*wave=*/0, /*vmcnt=*/-1, /*lgkmcnt=*/0); + b.checkSgprWrite(/*wave=*/0, /*reg=*/4); + EXPECT_FALSE(b.hasRace()); +} + +TEST(RaceDetector, TtmpWaw_WaitClearsPendingLoad) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.ttmpLoad(/*wave=*/0, /*ttmpBase=*/4, /*numRegs=*/1); + b.waitcnt(/*wave=*/0, /*vmcnt=*/-1, /*lgkmcnt=*/0); + b.checkTtmpWrite(/*wave=*/0, /*reg=*/4); + EXPECT_FALSE(b.hasRace()); +} + +TEST(RaceDetector, SgprWaw_CleanRegisterNoRace) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/4, /*numRegs=*/1); + b.checkSgprWrite(/*wave=*/0, /*reg=*/5); + EXPECT_FALSE(b.hasRace()); +} + +TEST(RaceDetector, SgprWaw_WideLoadReportsOnlyOverlappingRegister) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/3, /*numRegs=*/2); + b.checkSgprWrite(/*wave=*/0, /*reg=*/4); + + ASSERT_EQ(b.raceCount(), 1); + EXPECT_EQ(b.violations().front().index, 4); + EXPECT_EQ(b.violations().front().conflictingEvent, EventId{0}); + + b.clearViolations(); + b.checkSgprWrite(/*wave=*/0, /*reg=*/5); + EXPECT_FALSE(b.hasRace()); +} + +TEST(RaceDetector, Sgpr_WideReadReportsPendingEventOnce) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/3, /*numRegs=*/2); + b.checkSgprRead(/*wave=*/0, /*reg=*/3, /*width=*/2); + + ASSERT_EQ(b.raceCount(), 1); + const auto &violation = b.violations().front(); + EXPECT_EQ(violation.index, 3); + EXPECT_FALSE(violation.isWrite); + EXPECT_EQ(violation.conflictingEvent, EventId{0}); +} + +TEST(RaceDetector, SgprWaw_WideInstructionWriteReportsPendingEventOnce) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/3, /*numRegs=*/2); + b.checkSgprWrite(/*wave=*/0, /*reg=*/3, /*width=*/2); + + ASSERT_EQ(b.raceCount(), 1); + EXPECT_EQ(b.violations().front().index, 3); + EXPECT_EQ(b.violations().front().conflictingEvent, EventId{0}); +} + +TEST(RaceDetector, TtmpWaw_WideScalarLoadReportsPendingEventOnce) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.ttmpLoad(/*wave=*/0, /*ttmpBase=*/3, /*numRegs=*/2); + b.ttmpLoad(/*wave=*/0, /*ttmpBase=*/3, /*numRegs=*/2); + + ASSERT_EQ(b.raceCount(), 1); + EXPECT_EQ(b.violations().front().space, RaceViolation::Space::TTMP); + EXPECT_EQ(b.violations().front().index, 3); + EXPECT_EQ(b.violations().front().conflictingEvent, EventId{0}); +} + +TEST(RaceDetector, SgprWaw_PreservesEachPendingConflictingEvent) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/4, /*numRegs=*/1); + b.scalarLoad(/*wave=*/0, /*sgprBase=*/4, /*numRegs=*/1); + b.clearViolations(); + + b.checkSgprWrite(/*wave=*/0, /*reg=*/4); + + ASSERT_EQ(b.raceCount(), 2); + EXPECT_EQ(b.violations()[0].conflictingEvent, EventId{0}); + EXPECT_EQ(b.violations()[1].conflictingEvent, EventId{1}); +} + // ---- LDS cross-wave races ---- TEST(RaceDetector, LdsCrossWave_ZeroLdsSize) { diff --git a/emulation/rocjitsu/tests/race-detector/race_test_builder.h b/emulation/rocjitsu/tests/race-detector/race_test_builder.h index ab91dd68b2f1..8d94b76abed4 100644 --- a/emulation/rocjitsu/tests/race-detector/race_test_builder.h +++ b/emulation/rocjitsu/tests/race-detector/race_test_builder.h @@ -160,14 +160,25 @@ class RaceTestBuilder { waves_[wave]->checkVgprWriteLanes(reg, laneMask, byteMask); } - void checkSgprRead(int wave, int reg) { - waves_[wave]->checkScalarRead(RegisterRef{RegClass::SGPR, static_cast(reg), 1}); + void checkSgprRead(int wave, int reg, int width = 1) { + waves_[wave]->checkScalarRead( + RegisterRef{RegClass::SGPR, static_cast(reg), static_cast(width)}); + } + + void checkSgprWrite(int wave, int reg, int width = 1) { + waves_[wave]->checkScalarWrite( + RegisterRef{RegClass::SGPR, static_cast(reg), static_cast(width)}); } void checkTtmpRead(int wave, int reg) { waves_[wave]->checkScalarRead(RegisterRef{RegClass::TTMP, static_cast(reg), 1}); } + void checkTtmpWrite(int wave, int reg, int width = 1) { + waves_[wave]->checkScalarWrite( + RegisterRef{RegClass::TTMP, static_cast(reg), static_cast(width)}); + } + void checkLdsRead(int wave, int lane, int addr, int bytes) { detector_->validateRead(addr, WaveId{wave}, lane, bytes); }