intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2026-01-04 15:53:45 +08:00
Code Issues Packages Projects Releases Wiki Activity

fix: decanonize prefetch gpu va

Browse Source 
Signed-off-by: Bartosz Dunajski <bartosz.dunajski@intel.com>
This commit is contained in:
Bartosz Dunajski
2025-10-06 09:52:24 +00:00
committed by Compute-Runtime-Automation
parent 9ab15e28c1
commit 42371ee7bd
5 changed files with 26 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_xehp_and_later.cpp
View File

@@ -3028,7 +3028,9 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenNotEnoughIohSpaceWhenLaunchingKern
auto statePrefetch = genCmdCast<STATE_PREFETCH *>(*prefetchList);
ASSERT_NE(nullptr, statePrefetch);
EXPECT_EQ(ioh->getGraphicsAllocation()->getGpuAddress(), statePrefetch->getAddress());
auto gmmHelper = device->getNEODevice()->getGmmHelper();
EXPECT_EQ(gmmHelper->decanonize(ioh->getGraphicsAllocation()->getGpuAddress()), statePrefetch->getAddress());
}
HWTEST2_F(CommandListAppendLaunchKernel, givenDebugVariableWhenPrefetchingIsaThenLimitItsSize, IsAtLeastXeHpcCore) {

6
level_zero/core/test/unit_tests/sources/cmdlist/test_in_order_cmdlist_1.cpp
View File

@@ -2799,10 +2799,12 @@ HWTEST2_F(InOrderCmdListTests, givenIoqAndPrefetchEnabledWhenKernelIsAppendedThe
auto immCmdList = createImmCmdList<FamilyType::gfxCoreFamily>();
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
auto gmmHelper = device->getNEODevice()->getGmmHelper();
auto isaAllocation = kernel->getIsaAllocation();
auto isaAddress = isaAllocation->getGpuAddress() + kernel->getIsaOffsetInParentAllocation();
auto isaAddress = gmmHelper->decanonize(isaAllocation->getGpuAddress()) + kernel->getIsaOffsetInParentAllocation();
auto heap = immCmdList->getCmdContainer().getIndirectHeap(NEO::IndirectHeapType::indirectObject);
auto heapAddress = heap->getGraphicsAllocation()->getGpuAddress() + heap->getUsed();
auto heapAddress = gmmHelper->decanonize(heap->getGraphicsAllocation()->getGpuAddress()) + heap->getUsed();
zeCommandListAppendLaunchKernel(immCmdList->toHandle(), kernel->toHandle(), &groupCount, nullptr, 0, nullptr);

13
level_zero/core/test/unit_tests/sources/mutable_cmdlist/tests/mutable_cmdlist_kernels_tests.cpp
View File

@@ -6,6 +6,7 @@
*/
#include "shared/source/command_container/command_encoder.h"
#include "shared/source/gmm_helper/gmm_helper.h"
#include "shared/source/helpers/kernel_helpers.h"
#include "shared/source/indirect_heap/indirect_heap.h"
#include "shared/test/common/helpers/unit_test_helper.h"
@@ -573,12 +574,14 @@ HWTEST2_F(MutableCommandListKernelTest,
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::Parse::parseCommandBuffer(cmdList, prefetchCmdToPatch.pDestination, prefetchCmdToPatch.patchSize));
auto gmmHelper = device->getNEODevice()->getGmmHelper();
auto itor = cmdList.begin();
auto prefetchCmd = genCmdCast<STATE_PREFETCH *>(*itor);
ASSERT_NE(nullptr, prefetchCmd);
EXPECT_EQ(expectedIohPrefetchSize / MemoryConstants::cacheLineSize, prefetchCmd->getPrefetchSize());
EXPECT_EQ(mutation.kernelGroup->getIohForPrefetch()->getGpuAddress() + prefetchCmdToPatch.offset, prefetchCmd->getAddress());
EXPECT_EQ(gmmHelper->decanonize(mutation.kernelGroup->getIohForPrefetch()->getGpuAddress()) + prefetchCmdToPatch.offset, prefetchCmd->getAddress());
itor++;
for (size_t i = 0; i < expectedIohPrefetchPadding; i += sizeof(MI_NOOP)) {
@@ -590,7 +593,7 @@ HWTEST2_F(MutableCommandListKernelTest,
auto isaPrefetchCmd = genCmdCast<STATE_PREFETCH *>(*itor);
ASSERT_NE(nullptr, isaPrefetchCmd);
EXPECT_EQ(expectedIsaPrefetchSize / MemoryConstants::cacheLineSize, isaPrefetchCmd->getPrefetchSize());
EXPECT_EQ(kernel->getIsaAllocation()->getGpuAddress(), isaPrefetchCmd->getAddress());
EXPECT_EQ(gmmHelper->decanonize(kernel->getIsaAllocation()->getGpuAddress()), isaPrefetchCmd->getAddress());
itor++;
for (size_t i = 0; i < expectedIsaPrefetchPadding; i += sizeof(MI_NOOP)) {
@@ -673,10 +676,12 @@ HWTEST2_F(MutableCommandListKernelTest,
auto itor = cmdList.begin();
auto gmmHelper = device->getNEODevice()->getGmmHelper();
auto prefetchCmd = genCmdCast<STATE_PREFETCH *>(*itor);
ASSERT_NE(nullptr, prefetchCmd);
EXPECT_EQ(expectedIohPrefetchSize / MemoryConstants::cacheLineSize, prefetchCmd->getPrefetchSize());
EXPECT_EQ(mutation.kernelGroup->getIohForPrefetch()->getGpuAddress() + prefetchCmdToPatch.offset, prefetchCmd->getAddress());
EXPECT_EQ(gmmHelper->decanonize(mutation.kernelGroup->getIohForPrefetch()->getGpuAddress()) + prefetchCmdToPatch.offset, prefetchCmd->getAddress());
itor++;
for (size_t i = 0; i < expectedIohPrefetchPadding; i += sizeof(MI_NOOP)) {
@@ -688,7 +693,7 @@ HWTEST2_F(MutableCommandListKernelTest,
auto isaPrefetchCmd = genCmdCast<STATE_PREFETCH *>(*itor);
ASSERT_NE(nullptr, isaPrefetchCmd);
EXPECT_EQ(expectedIsaPrefetchSize / MemoryConstants::cacheLineSize, isaPrefetchCmd->getPrefetchSize());
EXPECT_EQ(kernel2->getIsaAllocation()->getGpuAddress(), isaPrefetchCmd->getAddress());
EXPECT_EQ(gmmHelper->decanonize(kernel2->getIsaAllocation()->getGpuAddress()), isaPrefetchCmd->getAddress());
itor++;
for (size_t i = 0; i < expectedIsaPrefetchPadding; i += sizeof(MI_NOOP)) {