Add test kernel for indirect allocations

Add kernel for testing indirect allocation
address getting passed by value.

Needed for future commit.

Related-To: NEO-6597

Signed-off-by: Dominik Dabek <dominik.dabek@intel.com>
This commit is contained in:
Dominik Dabek 2022-08-11 11:22:57 +02:00 committed by Compute-Runtime-Automation
parent bc605f413c
commit 0263ec0d84
4 changed files with 10 additions and 4 deletions

View File

@ -35,6 +35,12 @@ kernel void test(const global float *a, const global float *b,
printf("local_id = %d, global_id = %d \n", local_id, global_id);
}
__kernel void test_pointer_by_value(long address) {
if (address) {
((__global int *)address)[get_global_id(0)] += 1;
}
}
__kernel void test_get_global_sizes(__global uint *outGlobalSize) {
outGlobalSize[0] = get_global_size(0);
outGlobalSize[1] = get_global_size(1);

View File

@ -105,7 +105,7 @@ HWTEST_F(ModuleTest, givenBuiltinModuleWhenCreatedThenCorrectAllocationTypeIsUse
EXPECT_EQ(NEO::AllocationType::KERNEL_ISA_INTERNAL, kernel->getIsaAllocation()->getAllocationType());
}
HWTEST_F(ModuleTest, givenBlitterAvailableWhenCopyingPatchedSegementsThenIsaIsTransferredToAllocationWithBlitter) {
HWTEST_F(ModuleTest, givenBlitterAvailableWhenCopyingPatchedSegmentsThenIsaIsTransferredToAllocationWithBlitter) {
NEO::MockCompilerEnableGuard mock(true);
auto hwInfo = *NEO::defaultHwInfo;
hwInfo.featureTable.flags.ftrLocalMemory = true;
@ -157,7 +157,7 @@ HWTEST_F(ModuleTest, givenBlitterAvailableWhenCopyingPatchedSegementsThenIsaIsTr
const auto &hwInfoConfig = *HwInfoConfig::get(hwInfo.platform.eProductFamily);
if (hwInfoConfig.isBlitCopyRequiredForLocalMemory(hwInfo, *module->getKernelImmutableDataVector()[0]->getIsaGraphicsAllocation())) {
EXPECT_EQ(7u, blitterCalled);
EXPECT_EQ(8u, blitterCalled);
} else {
EXPECT_EQ(0u, blitterCalled);
}

View File

@ -38,7 +38,7 @@ components:
dest_dir: kernels_bin
type: git
branch: kernels_bin
revision: 1941-252
revision: 1941-263
kmdaf:
branch: kmdaf
dest_dir: kmdaf

View File

@ -901,7 +901,7 @@ TEST(LinkerInputTests, GivenGlobalAndLocalElfSymbolsWhenDecodingThenOnlyGlobalSy
EXPECT_TRUE(linkerInput.getTraits().exportsFunctions);
}
TEST(LinkerInputTests, GivenGlobalFunctionsInTwoSegementsWhenDecodingThenUnrecoverableIsCalled) {
TEST(LinkerInputTests, GivenGlobalFunctionsInTwoSegmentsWhenDecodingThenUnrecoverableIsCalled) {
NEO::LinkerInput linkerInput = {};
NEO::Elf::ElfFileHeader<NEO::Elf::EI_CLASS_64> header;
MockElf<NEO::Elf::EI_CLASS_64> elf64;