fix: remove not needed test files

merge kernel files to simple_kernels.cl

Signed-off-by: Mateusz Jablonski <mateusz.jablonski@intel.com>

opencl/test/unit_test/aub_tests/command_queue/enqueue_printf_kernel_aub_tests.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (C) 2020-2023 Intel Corporation
+ * Copyright (C) 2020-2024 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
@@ -33,7 +33,7 @@ class AUBPrintfKernelFixture : public AUBFixture,
         HelloWorldKernelFixture::tearDown();
         AUBFixture::tearDown();
     }
-    const char *programFile = "printf";
+    const char *programFile = "simple_kernels";
     const char *kernelName = "test_printf_number";
 };
 
@@ -57,4 +57,4 @@ HWTEST_F(AUBPrintfKernelFixture, GivenPrintfKernelThenEnqueuingSucceeds) {
 
     pCmdQ->enqueueKernel(pKernel, 1, offset, gws, lws, 0, 0, 0);
     pCmdQ->finish();
-}
+}

opencl/test/unit_test/context/driver_diagnostics_tests.h

@@ -265,10 +265,10 @@ struct PerformanceHintEnqueueKernelPrintfTest : public PerformanceHintEnqueueTes
 
     void SetUp() override {
         PerformanceHintEnqueueTest::SetUp();
-        createProgramFromBinary(context, context->getDevices(), "printf");
+        createProgramFromBinary(context, context->getDevices(), "simple_kernels");
         retVal = pProgram->build(pProgram->getDevices(), nullptr);
         ASSERT_EQ(CL_SUCCESS, retVal);
-        kernel = static_cast<KernelWhitebox *>(Kernel::create(pProgram, pProgram->getKernelInfoForKernel("test"), *context->getDevice(0), retVal));
+        kernel = static_cast<KernelWhitebox *>(Kernel::create(pProgram, pProgram->getKernelInfoForKernel("test_printf"), *context->getDevice(0), retVal));
         kernel->initializeLocalIdsCache();
 
         globalWorkGroupSize[0] = globalWorkGroupSize[1] = globalWorkGroupSize[2] = 1;

opencl/test/unit_test/fixtures/simple_arg_kernel_fixture.h

@@ -29,46 +29,6 @@ namespace NEO {
 class Kernel;
 class Program;
 
-template <typename T>
-inline const char *typeName(T &) {
-    return "unknown";
-}
-
-template <>
-inline const char *typeName(char &) {
-    return "char";
-}
-
-template <>
-inline const char *typeName(int &) {
-    return "int";
-}
-
-template <>
-inline const char *typeName(float &) {
-    return "float";
-}
-
-template <>
-inline const char *typeName(short &) {
-    return "short";
-}
-
-template <>
-inline const char *typeName(unsigned char &) {
-    return "unsigned char";
-}
-
-template <>
-inline const char *typeName(unsigned int &) {
-    return "unsigned int";
-}
-
-template <>
-inline const char *typeName(unsigned short &) {
-    return "unsigned short";
-}
-
 class SimpleArgKernelFixture : public ProgramFixture {
 
   public:
@@ -78,17 +38,6 @@ class SimpleArgKernelFixture : public ProgramFixture {
     void setUp(ClDevice *pDevice) {
         ProgramFixture::setUp();
 
-        std::string testFile;
-        int forTheName = 0;
-
-        testFile.append("simple_arg_");
-        testFile.append(typeName(forTheName));
-
-        auto pos = testFile.find(" ");
-        if (pos != (size_t)-1) {
-            testFile.replace(pos, 1, "_");
-        }
-
         auto deviceVector = toClDeviceVector(*pDevice);
         pContext = Context::create<MockContext>(nullptr, deviceVector, nullptr, nullptr, retVal);
         ASSERT_EQ(CL_SUCCESS, retVal);
@@ -97,7 +46,7 @@ class SimpleArgKernelFixture : public ProgramFixture {
         createProgramFromBinary(
             pContext,
             deviceVector,
-            testFile);
+            "simple_kernels");
         ASSERT_NE(nullptr, pProgram);
 
         retVal = pProgram->build(
@@ -108,7 +57,7 @@ class SimpleArgKernelFixture : public ProgramFixture {
         // create a kernel
         pKernel = Kernel::create<MockKernel>(
             pProgram,
-            pProgram->getKernelInfoForKernel("SimpleArg"),
+            pProgram->getKernelInfoForKernel("simple_arg_int"),
             *pDevice,
             retVal);
 

opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp

@@ -3255,16 +3255,20 @@ TEST(OfflineCompilerTest, givenUseLlvmBcFlagWhenBuildingIrBinaryThenProperTransl
 }
 
 TEST(OfflineCompilerTest, givenBinaryInputThenDontTruncateSourceAtFirstZero) {
-    std::vector<std::string> argvLlvm = {"ocloc", "-llvm_input", "-file", clFiles + "binary_with_zeroes", "-qq",
+    std::vector<std::string> argvLlvm = {"ocloc", "-llvm_input", "-file", "binary_with_zeroes", "-qq",
                                          "-device", gEnvironment->devicePrefix.c_str()};
     auto mockOfflineCompiler = std::make_unique<MockOfflineCompiler>();
+    uint8_t binaryWithZeros[64]{};
+    Source source{binaryWithZeros, sizeof(binaryWithZeros), "binary_with_zeroes"};
+    static_cast<MockOclocArgHelper *>(mockOfflineCompiler->argHelper)->inputs.push_back(source);
     mockOfflineCompiler->initialize(argvLlvm.size(), argvLlvm);
     mockOfflineCompiler->build();
     EXPECT_LT(0U, mockOfflineCompiler->sourceCode.size());
 
-    std::vector<std::string> argvSpirV = {"ocloc", "-spirv_input", "-file", clFiles + "binary_with_zeroes", "-qq",
+    std::vector<std::string> argvSpirV = {"ocloc", "-spirv_input", "-file", "binary_with_zeroes", "-qq",
                                           "-device", gEnvironment->devicePrefix.c_str()};
     mockOfflineCompiler = std::make_unique<MockOfflineCompiler>();
+    static_cast<MockOclocArgHelper *>(mockOfflineCompiler->argHelper)->inputs.push_back(source);
     mockOfflineCompiler->initialize(argvSpirV.size(), argvSpirV);
     mockOfflineCompiler->build();
     EXPECT_LT(0U, mockOfflineCompiler->sourceCode.size());
@@ -3906,27 +3910,6 @@ __kernel void shouldfail(global ushort *dst) {
     }
 }
 
-TEST(OfflineCompilerTest, givenInputOptionsFileWithSpecialCharsWhenOfflineCompilerIsInitializedThenCorrectOptionsAreSet) {
-    auto mockOfflineCompiler = std::unique_ptr<MockOfflineCompiler>(new MockOfflineCompiler());
-    ASSERT_NE(nullptr, mockOfflineCompiler);
-
-    ASSERT_TRUE(fileExists(clFiles + "simple_kernels_opts_options.txt"));
-
-    std::vector<std::string> argv = {
-        "ocloc",
-        "-q",
-        "-file",
-        clFiles + "simple_kernels_opts.cl",
-        "-device",
-        gEnvironment->devicePrefix.c_str()};
-
-    int retVal = mockOfflineCompiler->initialize(argv.size(), argv);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-
-    auto &options = mockOfflineCompiler->options;
-    EXPECT_STREQ(options.c_str(), "-cl-opt-disable -DDEF_WAS_SPECIFIED=1 -DARGS=\", const __global int *arg1, float arg2, const __global int *arg3, float arg4\"");
-}
-
 TEST(OfflineCompilerTest, givenInputOptionsAndOclockOptionsFileWithForceStosOptWhenOfflineCompilerIsInitializedThenCompilerOptionGreaterThan4gbBuffersRequiredIsNotApplied) {
     auto mockOfflineCompiler = std::unique_ptr<MockOfflineCompiler>(new MockOfflineCompiler());
     ASSERT_NE(nullptr, mockOfflineCompiler);
@@ -4198,7 +4181,7 @@ TEST(OfflineCompilerTest, givenDeviceSpecificKernelFileWhenCompilerIsInitialized
  *
  */
 
--cl-opt-disable
+-cl-opt-disable -DDEF_WAS_SPECIFIED=1 -DARGS=", const __global int *arg1, float arg2, const __global int *arg3, float arg4"
 )===";
     Source optionsSource{reinterpret_cast<const uint8_t *>(options), sizeof(options), "emptykernel_options.txt"};
     static_cast<MockOclocArgHelper *>(mockOfflineCompiler->argHelper)->inputs.push_back(optionsSource);
@@ -4213,7 +4196,7 @@ TEST(OfflineCompilerTest, givenDeviceSpecificKernelFileWhenCompilerIsInitialized
 
     int retVal = mockOfflineCompiler->initialize(argv.size(), argv);
     EXPECT_EQ(OCLOC_SUCCESS, retVal);
-    EXPECT_STREQ("-cl-opt-disable", mockOfflineCompiler->options.c_str());
+    EXPECT_STREQ("-cl-opt-disable -DDEF_WAS_SPECIFIED=1 -DARGS=\", const __global int *arg1, float arg2, const __global int *arg3, float arg4\"", mockOfflineCompiler->options.c_str());
 }
 
 TEST(OfflineCompilerTest, givenHexadecimalRevisionIdWhenCompilerIsInitializedThenPassItToHwInfo) {
@@ -4496,11 +4479,13 @@ TEST(OclocCompile, givenPackedDeviceBinaryFormatWhenGeneratingElfBinaryThenItIsR
 TEST(OclocCompile, givenSpirvInputThenDontGenerateSpirvFile) {
     MockOfflineCompiler ocloc;
 
+    Source source{reinterpret_cast<const uint8_t *>(spirvMagic.data()), spirvMagic.size(), "some_file.spv"};
+    static_cast<MockOclocArgHelper *>(ocloc.argHelper)->inputs.push_back(source);
     std::vector<std::string> argv = {
         "ocloc",
         "-q",
         "-file",
-        clFiles + "binary_with_zeroes",
+        "some_file.spv",
         "-out_dir",
         "offline_compiler_test",
         "-device",
@@ -4511,9 +4496,9 @@ TEST(OclocCompile, givenSpirvInputThenDontGenerateSpirvFile) {
     ASSERT_EQ(0, retVal);
     retVal = ocloc.build();
     EXPECT_EQ(0, retVal);
-    EXPECT_FALSE(compilerOutputExists("offline_compiler_test/binary_with_zeroes", "gen"));
-    EXPECT_TRUE(compilerOutputExists("offline_compiler_test/binary_with_zeroes", "bin"));
-    EXPECT_FALSE(compilerOutputExists("offline_compiler_test/binary_with_zeroes", "spv"));
+    EXPECT_FALSE(compilerOutputExists("offline_compiler_test/some_file", "gen"));
+    EXPECT_TRUE(compilerOutputExists("offline_compiler_test/some_file", "bin"));
+    EXPECT_FALSE(compilerOutputExists("offline_compiler_test/some_file", "spv"));
 }
 
 TEST(OclocCompile, givenFormatFlagWithKnownFormatPassedThenEnforceSpecifiedFormatAccordingly) {

opencl/test/unit_test/program/program_nonuniform.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (C) 2018-2023 Intel Corporation
+ * Copyright (C) 2018-2024 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
@@ -178,7 +178,7 @@ class ProgramNonUniformTest : public ContextFixture,
 TEST_F(ProgramNonUniformTest, GivenCl21WhenExecutingKernelWithNonUniformThenEnqueueSucceeds) {
     REQUIRE_OCL_21_OR_SKIP(defaultHwInfo);
 
-    createProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
+    createProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
     auto mockProgram = pProgram;
     ASSERT_NE(nullptr, mockProgram);
 
@@ -219,7 +219,7 @@ TEST_F(ProgramNonUniformTest, GivenCl21WhenExecutingKernelWithNonUniformThenEnqu
 TEST_F(ProgramNonUniformTest, GivenCl20WhenExecutingKernelWithNonUniformThenEnqueueSucceeds) {
     REQUIRE_OCL_21_OR_SKIP(defaultHwInfo);
 
-    createProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
+    createProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
     auto mockProgram = pProgram;
     ASSERT_NE(nullptr, mockProgram);
 
@@ -258,7 +258,7 @@ TEST_F(ProgramNonUniformTest, GivenCl20WhenExecutingKernelWithNonUniformThenEnqu
 }
 
 TEST_F(ProgramNonUniformTest, GivenCl12WhenExecutingKernelWithNonUniformThenInvalidWorkGroupSizeIsReturned) {
-    createProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
+    createProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
     auto mockProgram = pProgram;
     ASSERT_NE(nullptr, mockProgram);
 

opencl/test/unit_test/program/program_tests.cpp

@@ -1427,7 +1427,7 @@ class PatchTokenFromBinaryTest : public ProgramSimpleFixture {
 using PatchTokenTests = Test<PatchTokenFromBinaryTest>;
 
 TEST_F(PatchTokenTests, WhenBuildingProgramThenGwsIsSet) {
-    createProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
+    createProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
 
     ASSERT_NE(nullptr, pProgram);
     retVal = pProgram->build(
@@ -1436,7 +1436,7 @@ TEST_F(PatchTokenTests, WhenBuildingProgramThenGwsIsSet) {
 
     ASSERT_EQ(CL_SUCCESS, retVal);
 
-    auto pKernelInfo = pProgram->getKernelInfo("test", rootDeviceIndex);
+    auto pKernelInfo = pProgram->getKernelInfo("test_get_global_size", rootDeviceIndex);
 
     ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.globalWorkSize[0]);
     ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.globalWorkSize[1]);
@@ -1446,7 +1446,7 @@ TEST_F(PatchTokenTests, WhenBuildingProgramThenGwsIsSet) {
 TEST_F(PatchTokenTests, WhenBuildingProgramThenConstantKernelArgsAreAvailable) {
     // PATCH_TOKEN_STATELESS_CONSTANT_MEMORY_OBJECT_KERNEL_ARGUMENT
 
-    createProgramFromBinary(pContext, pContext->getDevices(), "test_basic_constant");
+    createProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
 
     ASSERT_NE(nullptr, pProgram);
     retVal = pProgram->build(
@@ -3059,7 +3059,7 @@ TEST_F(ProgramBinTest, givenPrintProgramBinaryProcessingTimeSetWhenBuildProgramT
     debugManager.flags.PrintProgramBinaryProcessingTime.set(true);
     testing::internal::CaptureStdout();
 
-    createProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
+    createProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
 
     auto retVal = pProgram->build(
         pProgram->getDevices(),

opencl/test/unit_test/test_files/kernel_data_param.cl

@@ -1,20 +0,0 @@
-/*
- * Copyright (C) 2018-2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
-__kernel void test(__global int *dst) {
-    int tid = get_global_id(0);
-    int n = get_global_size(0);
-
-    dst[tid] = n;
-};
-
-__kernel void test_get_local_size(__global int *dst) {
-    int tid = get_global_id(0);
-    int n = get_local_size(0);
-
-    dst[tid] = n;
-};

opencl/test/unit_test/test_files/printf.cl

@@ -1,14 +0,0 @@
-/*
- * Copyright (C) 2018-2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
-__kernel void test() {
-    printf("OpenCL\n");
-}
-
-__kernel void test_printf_number(__global uint* in) {
-    printf("in[0] = %d\n", in[0]);
-}

opencl/test/unit_test/test_files/simple_arg_int.cl

@@ -1,11 +0,0 @@
-/*
- * Copyright (C) 2018-2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
-__kernel void SimpleArg(int src, __global int *dst) {
-    int id = (int)get_global_id(0);
-    dst[id] = src;
-}

opencl/test/unit_test/test_files/simple_kernels.cl

@@ -1,5 +1,5 @@
 /*
- * Copyright (C) 2018-2021 Intel Corporation
+ * Copyright (C) 2018-2024 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
@@ -46,7 +46,7 @@ __kernel void simple_kernel_4() {
 }
 
 __kernel void simple_kernel_5(__global uint *dst) {
-    //first uint holds the total work item count
+    // first uint holds the total work item count
     atomic_inc(dst);
     uint groupIdX = get_group_id(0);
     uint groupIdY = get_group_id(1);
@@ -57,7 +57,7 @@ __kernel void simple_kernel_5(__global uint *dst) {
     uint groupCountZ = get_num_groups(2);
 
     __global uint *groupCounters = dst + 1;
-    //store current group position in 3D array
+    // store current group position in 3D array
     uint destination = groupIdZ * groupCountY * groupCountX + groupIdY * groupCountX + groupIdX;
     atomic_inc(&groupCounters[destination]);
 }
@@ -89,7 +89,7 @@ __kernel void simple_kernel_6(__global uint *dst, __constant uint2 *src, uint sc
 
 typedef long16 TYPE;
 __attribute__((reqd_work_group_size(32, 1, 1))) // force LWS to 32
-__attribute__((intel_reqd_sub_group_size(16)))   // force SIMD to 16
+__attribute__((intel_reqd_sub_group_size(16)))  // force SIMD to 16
 __kernel void
 simple_kernel_7(__global int *resIdx, global TYPE *src, global TYPE *dst) {
     size_t lid = get_local_id(0);
@@ -128,3 +128,37 @@ __kernel void simple_kernel_9(__global uint *dst) {
     uint offset = get_max_sub_group_size() * get_sub_group_id();
     dst[get_sub_group_local_id() + offset] = get_local_id(0);
 }
+
+__kernel void constant_kernel(__global float *out, __constant float *tmpF, __constant int *tmpI) {
+    int tid = get_global_id(0);
+
+    float ftmp = tmpF[tid];
+    float Itmp = tmpI[tid];
+    out[tid] = ftmp * Itmp;
+}
+
+__kernel void test_get_global_size(__global int *dst) {
+    int tid = get_global_id(0);
+    int n = get_global_size(0);
+
+    dst[tid] = n;
+};
+
+__kernel void test_get_local_size(__global int *dst) {
+    int tid = get_global_id(0);
+    int n = get_local_size(0);
+
+    dst[tid] = n;
+};
+
+__kernel void test_printf() {
+    printf("OpenCL\n");
+}
+
+__kernel void test_printf_number(__global uint *in) {
+    printf("in[0] = %d\n", in[0]);
+}
+__kernel void simple_arg_int(int src, __global int *dst) {
+    int id = (int)get_global_id(0);
+    dst[id] = src;
+}

opencl/test/unit_test/test_files/simple_kernels_opts.cl

@@ -1,18 +0,0 @@
-/*
- * Copyright (C) 2018-2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
-__kernel void SimpleArg(int src, __global int *dst) {
-    int id = (int)get_global_id(0);
-
-#ifdef DEF_WAS_SPECIFIED
-    int val = 1;
-#else
-    // fail to compile
-#endif
-
-    dst[id] = src + val;
-}

opencl/test/unit_test/test_files/simple_kernels_opts_options.txt

@@ -1,49 +0,0 @@
-/*
- * Copyright (C) 2018-2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
--cl-opt-disable -DDEF_WAS_SPECIFIED=1 -DARGS=", const __global int *arg1, float arg2, const __global int *arg3, float arg4"
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-

opencl/test/unit_test/test_files/test_basic_constant.cl

@@ -1,14 +0,0 @@
-/*
- * Copyright (C) 2018-2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
-__kernel void constant_kernel(__global float *out, __constant float *tmpF, __constant int *tmpI) {
-    int tid = get_global_id(0);
-
-    float ftmp = tmpF[tid];
-    float Itmp = tmpI[tid];
-    out[tid] = ftmp * Itmp;
-}