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 fixture with a simple stateless copy kernel

Browse Source 
Signed-off-by: Slawomir Milczarek <slawomir.milczarek@intel.com>
This commit is contained in:
Slawomir Milczarek
2020-12-09 10:54:19 +00:00
committed by Compute-Runtime-Automation
parent 408373412d
commit e7772143a2
6 changed files with 65 additions and 4 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

@ -265,6 +265,45 @@ class SimpleKernelStatelessFixture : public ProgramFixture {
cl_int retVal = CL_SUCCESS;
};
class StatelessCopyKernelFixture : 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),
"stateless_copy_buffer");
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("StatelessCopyBuffer"),
&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;

4
opencl/test/unit_test/program/kernel_data.cpp
View File

@ -326,7 +326,7 @@ TEST_F(KernelDataTest, GivenExecutionEnvironmentNoReqdWorkGroupSizeWhenBuildingT
EXPECT_EQ_VAL(0, pKernelInfo->kernelDescriptor.kernelAttributes.requiredWorkgroupSize[0]);
EXPECT_EQ_VAL(0, pKernelInfo->kernelDescriptor.kernelAttributes.requiredWorkgroupSize[1]);
EXPECT_EQ_VAL(0, pKernelInfo->kernelDescriptor.kernelAttributes.requiredWorkgroupSize[2]);
EXPECT_FALSE(pKernelInfo->hasStatelessAccessToHostMemory);
EXPECT_FALSE(pKernelInfo->hasIndirectStatelessAccess);
}
TEST_F(KernelDataTest, GivenExecutionEnvironmentWhenBuildingThenProgramIsCorrect) {
@ -364,7 +364,7 @@ TEST_F(KernelDataTest, GivenExecutionEnvironmentWhenBuildingThenProgramIsCorrect
EXPECT_EQ(16u, pKernelInfo->kernelDescriptor.kernelAttributes.requiredWorkgroupSize[1]);
EXPECT_EQ(8u, pKernelInfo->kernelDescriptor.kernelAttributes.requiredWorkgroupSize[2]);
EXPECT_TRUE(pKernelInfo->requiresSshForBuffers);
EXPECT_TRUE(pKernelInfo->hasStatelessAccessToHostMemory);
EXPECT_TRUE(pKernelInfo->hasIndirectStatelessAccess);
}
TEST_F(KernelDataTest, GivenExecutionEnvironmentCompiledForGreaterThan4gbBuffersWhenBuildingThenProgramIsCorrect) {

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

@ -0,0 +1,14 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
__kernel void StatelessCopyBuffer(
const __global uchar* src,
__global uchar* dst)
{
uint id = get_global_id(0);
dst[id] = src[id];
}

8
opencl/test/unit_test/test_files/stateless_copy_buffer_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