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

Add kernel fixture with stateless indirect access

Browse Source 
Signed-off-by: Slawomir Milczarek <slawomir.milczarek@intel.com>
This commit is contained in:
Slawomir Milczarek
2020-12-14 14:06:19 +00:00
committed by Compute-Runtime-Automation
parent 7bfa71636a
commit b2cae02f66
7 changed files with 88 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

39
opencl/test/unit_test/fixtures/simple_arg_kernel_fixture.h
View File

@ -304,6 +304,45 @@ class StatelessCopyKernelFixture : public ProgramFixture {
cl_int retVal = CL_SUCCESS;
};
class StatelessKernelWithIndirectAccessFixture : public ProgramFixture {
public:
DebugManagerStateRestore restorer;
using ProgramFixture::SetUp;
protected:
void SetUp(ClDevice *device, Context *context) {
ProgramFixture::SetUp();
DebugManager.flags.DisableStatelessToStatefulOptimization.set(true);
DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(false);
CreateProgramFromBinary(
context,
toClDeviceVector(*device),
"indirect_access_kernel");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
pProgram->getDevices(),
CompilerOptions::greaterThan4gbBuffersRequired.data(),
false);
ASSERT_EQ(CL_SUCCESS, retVal);
kernel.reset(Kernel::create<MockKernel>(
pProgram,
pProgram->getKernelInfosForKernel("testIndirect"),
&retVal));
ASSERT_NE(nullptr, kernel);
ASSERT_EQ(CL_SUCCESS, retVal);
}
void TearDown() override {
ProgramFixture::TearDown();
}
std::unique_ptr<Kernel> kernel = nullptr;
cl_int retVal = CL_SUCCESS;
};
class BindlessKernelFixture : public ProgramFixture {
public:
using ProgramFixture::SetUp;

14
opencl/test/unit_test/test_files/indirect_access_kernel.cl Normal file
View File

@ -0,0 +1,14 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
__kernel void testIndirect(__global long* buf) {
size_t gid = get_global_id(0);
if (gid == 0) {
char* val = (char*)buf[0];
*val = 1;
}
}

8
opencl/test/unit_test/test_files/indirect_access_kernel_internal_options.txt Normal file
View File

@ -0,0 +1,8 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
-cl-intel-greater-than-4GB-buffer-required

1
shared/source/memory_manager/CMakeLists.txt
View File

@ -51,6 +51,7 @@ set(NEO_CORE_MEMORY_MANAGER
${CMAKE_CURRENT_SOURCE_DIR}/surface.h
${CMAKE_CURRENT_SOURCE_DIR}/unified_memory_manager.cpp
${CMAKE_CURRENT_SOURCE_DIR}/unified_memory_manager.h
${CMAKE_CURRENT_SOURCE_DIR}/${BRANCH_DIR_SUFFIX}/unified_memory_manager_extra.cpp
)
set_property(GLOBAL PROPERTY NEO_CORE_MEMORY_MANAGER ${NEO_CORE_MEMORY_MANAGER})

9
shared/source/memory_manager/unified_memory_manager.cpp
View File

@ -166,14 +166,7 @@ void *SVMAllocsManager::createUnifiedMemoryAllocation(uint32_t rootDeviceIndex,
size_t alignedSize = alignUp<size_t>(size, MemoryConstants::pageSize64k);
GraphicsAllocation::AllocationType allocationType = GraphicsAllocation::AllocationType::BUFFER_HOST_MEMORY;
if (memoryProperties.memoryType == InternalMemoryType::DEVICE_UNIFIED_MEMORY) {
if (memoryProperties.allocationFlags.allocFlags.allocWriteCombined) {
allocationType = GraphicsAllocation::AllocationType::WRITE_COMBINED;
} else {
allocationType = GraphicsAllocation::AllocationType::BUFFER;
}
}
GraphicsAllocation::AllocationType allocationType = getGraphicsAllocationType(memoryProperties);
AllocationProperties unifiedMemoryProperties{rootDeviceIndex,
true,

1
shared/source/memory_manager/unified_memory_manager.h
View File

@ -139,6 +139,7 @@ class SVMAllocsManager {
protected:
void *createZeroCopySvmAllocation(uint32_t rootDeviceIndex, size_t size, const SvmAllocationProperties &svmProperties, const DeviceBitfield &deviceBitfield);
GraphicsAllocation::AllocationType getGraphicsAllocationType(const UnifiedMemoryProperties &unifiedMemoryProperties) const;
void freeZeroCopySvmAllocation(SvmAllocationData *svmData);

24
shared/source/memory_manager/unified_memory_manager_extra.cpp Normal file
View File

@ -0,0 +1,24 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "shared/source/memory_manager/unified_memory_manager.h"
namespace NEO {
GraphicsAllocation::AllocationType SVMAllocsManager::getGraphicsAllocationType(const UnifiedMemoryProperties &unifiedMemoryProperties) const {
GraphicsAllocation::AllocationType allocationType = GraphicsAllocation::AllocationType::BUFFER_HOST_MEMORY;
if (unifiedMemoryProperties.memoryType == InternalMemoryType::DEVICE_UNIFIED_MEMORY) {
if (unifiedMemoryProperties.allocationFlags.allocFlags.allocWriteCombined) {
allocationType = GraphicsAllocation::AllocationType::WRITE_COMBINED;
} else {
allocationType = GraphicsAllocation::AllocationType::BUFFER;
}
}
return allocationType;
}
} // namespace NEO