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

Revert "Use total SLM instead of inline in error checking"

Browse Source 
This reverts commit e1c49d0f23.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
This commit is contained in:
Compute-Runtime-Validation
2022-10-22 08:46:42 +02:00
committed by Compute-Runtime-Automation
parent 9ad3f6190f
commit 052164c559
4 changed files with 21 additions and 43 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
opencl/source/api/api.cpp
View File

@ -3511,10 +3511,10 @@ cl_int CL_API_CALL clEnqueueNDRangeKernel(cl_command_queue commandQueue,
Kernel *pKernel = pMultiDeviceKernel->getKernel(pCommandQueue->getDevice().getRootDeviceIndex());
auto localMemSize = static_cast<uint32_t>(pCommandQueue->getDevice().getDeviceInfo().localMemSize);
auto slmTotalSize = pKernel->getSlmTotalSize();
auto slmInlineSize = pKernel->getDescriptor().kernelAttributes.slmInlineSize;
if (slmTotalSize > 0 && localMemSize < slmTotalSize) {
PRINT_DEBUG_STRING(NEO::DebugManager.flags.PrintDebugMessages.get(), stderr, "Size of SLM (%u) larger than available (%u)\n", slmTotalSize, localMemSize);
if (slmInlineSize > 0 && localMemSize < slmInlineSize) {
PRINT_DEBUG_STRING(NEO::DebugManager.flags.PrintDebugMessages.get(), stderr, "Size of SLM (%u) larger than available (%u)\n", slmInlineSize, localMemSize);
retVal = CL_OUT_OF_RESOURCES;
TRACING_EXIT(ClEnqueueNdRangeKernel, &retVal);
return retVal;

12
opencl/source/kernel/kernel.h
View File

@ -101,14 +101,12 @@ class Kernel : public ReferenceTrackedObject<Kernel> {
pKernel = nullptr;
}
if (pKernel) {
auto localMemSize = static_cast<uint32_t>(clDevice.getDevice().getDeviceInfo().localMemSize);
auto slmTotalSize = pKernel->getSlmTotalSize();
auto localMemSize = static_cast<uint32_t>(clDevice.getDevice().getDeviceInfo().localMemSize);
auto slmInlineSize = kernelInfo.kernelDescriptor.kernelAttributes.slmInlineSize;
if (slmTotalSize > 0 && localMemSize < slmTotalSize) {
PRINT_DEBUG_STRING(NEO::DebugManager.flags.PrintDebugMessages.get(), stderr, "Size of SLM (%u) larger than available (%u)\n", slmTotalSize, localMemSize);
retVal = CL_OUT_OF_RESOURCES;
}
if (slmInlineSize > 0 && localMemSize < slmInlineSize) {
PRINT_DEBUG_STRING(NEO::DebugManager.flags.PrintDebugMessages.get(), stderr, "Size of SLM (%u) larger than available (%u)\n", slmInlineSize, localMemSize);
retVal = CL_OUT_OF_RESOURCES;
}
if (errcodeRet) {

6
opencl/test/unit_test/api/cl_enqueue_nd_range_kernel_tests.inl
View File

@ -55,7 +55,7 @@ TEST_F(clEnqueueNDRangeKernelTests, GivenKernelWithSlmSizeExceedingLocalMemorySi
auto localMemSize = static_cast<uint32_t>(pDevice->getDevice().getDeviceInfo().localMemSize);
pKernel->setTotalSLMSize(localMemSize - 10u);
pProgram->mockKernelInfo.kernelDescriptor.kernelAttributes.slmInlineSize = localMemSize - 10u;
retVal = clEnqueueNDRangeKernel(
pCommandQueue,
pMultiDeviceKernel,
@ -74,7 +74,7 @@ TEST_F(clEnqueueNDRangeKernelTests, GivenKernelWithSlmSizeExceedingLocalMemorySi
::testing::internal::CaptureStderr();
pKernel->setTotalSLMSize(localMemSize + 10u);
pProgram->mockKernelInfo.kernelDescriptor.kernelAttributes.slmInlineSize = localMemSize + 10u;
retVal = clEnqueueNDRangeKernel(
pCommandQueue,
pMultiDeviceKernel,
@ -89,7 +89,7 @@ TEST_F(clEnqueueNDRangeKernelTests, GivenKernelWithSlmSizeExceedingLocalMemorySi
EXPECT_EQ(CL_OUT_OF_RESOURCES, retVal);
output = testing::internal::GetCapturedStderr();
const auto &slmInlineSize = pKernel->getSlmTotalSize();
const auto &slmInlineSize = pProgram->mockKernelInfo.kernelDescriptor.kernelAttributes.slmInlineSize;
std::string expectedOutput = "Size of SLM (" + std::to_string(slmInlineSize) + ") larger than available (" + std::to_string(localMemSize) + ")\n";
EXPECT_EQ(expectedOutput, output);
}

40
opencl/test/unit_test/kernel/kernel_tests.cpp
View File

@ -2992,7 +2992,6 @@ class KernelCreateTest : public ::testing::Test {
struct MockKernel {
MockKernel(MockProgram *, const KernelInfo &, ClDevice &) {}
int initialize() { return -1; };
uint32_t getSlmTotalSize() const { return 0u; };
};
MockProgram mockProgram{};
@ -3007,27 +3006,6 @@ TEST_F(KernelCreateTest, whenInitFailedThenReturnNull) {
}
TEST_F(KernelCreateTest, whenSlmSizeExceedsLocalMemorySizeThenDebugMsgErrIsPrintedAndOutOfResourcesIsReturned) {
struct MockKernel {
MockKernel(MockProgram *, const KernelInfo &, ClDevice &clDevice) {
deviceLocalMemSize = static_cast<uint32_t>(clDevice.getDevice().getDeviceInfo().localMemSize);
}
int initialize() { return 0; };
uint32_t getSlmTotalSize() const {
return deviceLocalMemSize - 10u;
};
uint32_t deviceLocalMemSize = 0u;
};
struct MockKernelExceedSLM {
MockKernelExceedSLM(MockProgram *, const KernelInfo &, ClDevice &clDevice) {
deviceLocalMemSize = static_cast<uint32_t>(clDevice.getDevice().getDeviceInfo().localMemSize);
}
int initialize() { return 0; };
uint32_t getSlmTotalSize() const {
return deviceLocalMemSize + 10u;
};
uint32_t deviceLocalMemSize = 0u;
};
DebugManagerStateRestore dbgRestorer;
DebugManager.flags.PrintDebugMessages.set(true);
@ -3038,8 +3016,9 @@ TEST_F(KernelCreateTest, whenSlmSizeExceedsLocalMemorySizeThenDebugMsgErrIsPrint
auto localMemSize = static_cast<uint32_t>(mockProgram.mDevice.getDevice().getDeviceInfo().localMemSize);
std::unique_ptr<MockKernel> kernel0(Kernel::create<MockKernel>(&mockProgram, info, mockProgram.mDevice, &retVal));
EXPECT_NE(nullptr, kernel0.get());
info.kernelDescriptor.kernelAttributes.slmInlineSize = localMemSize - 10u;
auto ret = Kernel::create<MockKernel>(&mockProgram, info, mockProgram.mDevice, &retVal);
EXPECT_EQ(nullptr, ret);
EXPECT_NE(CL_OUT_OF_RESOURCES, retVal);
std::string output = testing::internal::GetCapturedStderr();
@ -3049,19 +3028,20 @@ TEST_F(KernelCreateTest, whenSlmSizeExceedsLocalMemorySizeThenDebugMsgErrIsPrint
retVal = 0;
std::unique_ptr<MockKernelExceedSLM> kernel1(Kernel::create<MockKernelExceedSLM>(&mockProgram, info, mockProgram.mDevice, &retVal));
EXPECT_NE(nullptr, kernel1.get());
info.kernelDescriptor.kernelAttributes.slmInlineSize = localMemSize + 10u;
ret = Kernel::create<MockKernel>(&mockProgram, info, mockProgram.mDevice, &retVal);
EXPECT_EQ(nullptr, ret);
EXPECT_EQ(CL_OUT_OF_RESOURCES, retVal);
output = testing::internal::GetCapturedStderr();
const auto &slmTotalSize = localMemSize + 10u;
std::string expectedOutput = "Size of SLM (" + std::to_string(slmTotalSize) + ") larger than available (" + std::to_string(localMemSize) + ")\n";
const auto &slmInlineSize = info.kernelDescriptor.kernelAttributes.slmInlineSize;
std::string expectedOutput = "Size of SLM (" + std::to_string(slmInlineSize) + ") larger than available (" + std::to_string(localMemSize) + ")\n";
EXPECT_EQ(expectedOutput, output);
::testing::internal::CaptureStderr();
std::unique_ptr<MockKernelExceedSLM> kernel2(Kernel::create<MockKernelExceedSLM>(&mockProgram, info, mockProgram.mDevice, nullptr));
EXPECT_NE(nullptr, kernel2.get());
ret = Kernel::create<MockKernel>(&mockProgram, info, mockProgram.mDevice, nullptr);
EXPECT_EQ(nullptr, ret);
output = testing::internal::GetCapturedStderr();
EXPECT_EQ(expectedOutput, output);