diff --git a/offline_compiler/offline_compiler.cpp b/offline_compiler/offline_compiler.cpp index 89f777a68c..b459d879a8 100644 --- a/offline_compiler/offline_compiler.cpp +++ b/offline_compiler/offline_compiler.cpp @@ -138,7 +138,8 @@ int OfflineCompiler::buildSourceCode() { if (!inputFileLlvm) { IGC::CodeType::CodeType_t intermediateRepresentation = useLlvmText ? IGC::CodeType::llvmLl : IGC::CodeType::llvmBc; - auto fclSrc = CIF::Builtins::CreateConstBuffer(fclMain.get(), sourceCode.c_str(), sourceCode.size()); + // sourceCode.size() returns the number of characters without null terminated char + auto fclSrc = CIF::Builtins::CreateConstBuffer(fclMain.get(), sourceCode.c_str(), sourceCode.size() + 1); auto fclOptions = CIF::Builtins::CreateConstBuffer(fclMain.get(), options.c_str(), options.size()); auto fclInternalOptions = CIF::Builtins::CreateConstBuffer(fclMain.get(), internalOptions.c_str(), internalOptions.size()); diff --git a/unit_tests/CMakeLists.txt b/unit_tests/CMakeLists.txt index 7e1fbceb50..dc7eb0186f 100644 --- a/unit_tests/CMakeLists.txt +++ b/unit_tests/CMakeLists.txt @@ -318,6 +318,7 @@ set(TEST_KERNELS test_files/CopyBuffer_simd32.cl test_files/CopyBuffer_simd8.cl test_files/copybuffer_with_header.cl + test_files/emptykernel.cl test_files/kernel_data_param.cl test_files/kernel_num_args.cl test_files/media_kernels_backend.cl diff --git a/unit_tests/offline_compiler/offline_compiler_tests.cpp b/unit_tests/offline_compiler/offline_compiler_tests.cpp index d2465e1218..a712343c6e 100644 --- a/unit_tests/offline_compiler/offline_compiler_tests.cpp +++ b/unit_tests/offline_compiler/offline_compiler_tests.cpp @@ -488,6 +488,27 @@ TEST(OfflineCompilerTest, buildSourceCode) { EXPECT_NE(0u, mockOfflineCompiler->getGenBinarySize()); } +TEST(OfflineCompilerTest, GivenKernelWhenNoCharAfterKernelSourceThenBuildWithSuccess) { + auto mockOfflineCompiler = std::unique_ptr(new MockOfflineCompiler()); + ASSERT_NE(nullptr, mockOfflineCompiler); + + auto retVal = mockOfflineCompiler->buildSourceCode(); + EXPECT_EQ(CL_INVALID_PROGRAM, retVal); + + const char *argv[] = { + "cloc", + "-file", + "test_files/emptykernel.cl", + "-device", + gEnvironment->devicePrefix.c_str()}; + + retVal = mockOfflineCompiler->initialize(ARRAY_COUNT(argv), argv); + EXPECT_EQ(CL_SUCCESS, retVal); + + retVal = mockOfflineCompiler->buildSourceCode(); + EXPECT_EQ(CL_SUCCESS, retVal); +} + TEST(OfflineCompilerTest, generateElfBinary) { auto mockOfflineCompiler = std::unique_ptr(new MockOfflineCompiler()); ASSERT_NE(nullptr, mockOfflineCompiler); diff --git a/unit_tests/test_files/emptykernel.cl b/unit_tests/test_files/emptykernel.cl new file mode 100644 index 0000000000..492135b1e7 --- /dev/null +++ b/unit_tests/test_files/emptykernel.cl @@ -0,0 +1,25 @@ +/* + * Copyright (c) 2018, Intel Corporation + * + * Permission is hereby granted, free of charge, to any person obtaining a + * copy of this software and associated documentation files (the "Software"), + * to deal in the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included + * in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS + * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR + * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, + * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR + * OTHER DEALINGS IN THE SOFTWARE. + */ + /* No character after the last "}" */ + __kernel void empty() +{ +} \ No newline at end of file