Do not force cl_khr_3d_image_writes extension when compiling kernels

Update usage of SUPPORTED_IMAGES flag and do not use images when disabled.
Use SUPPORTED_2_0 only on fully OCL 2.1 conformant platforms.

Signed-off-by: Filip Hazubski <filip.hazubski@intel.com>

shared/offline_compiler/source/offline_compiler.cpp

@@ -361,6 +361,26 @@ int OfflineCompiler::initialize(size_t numArgs, const std::vector<std::string> &
         }
     }
 
+    retVal = deviceName.empty() ? 0 : getHardwareInfo(deviceName.c_str());
+    if (retVal != 0) {
+        argHelper->printf("Error: Cannot get HW Info for device %s.\n", deviceName.c_str());
+        return retVal;
+    }
+    if (deviceName.empty()) {
+        CompilerOptions::concatenateAppend(internalOptions, "-cl-ext=-all,+cl_khr_3d_image_writes");
+    } else {
+        auto oclVersion = getOclVersionCompilerInternalOption(hwInfo.capabilityTable.clVersionSupport);
+        CompilerOptions::concatenateAppend(internalOptions, oclVersion);
+
+        std::string extensionsList = getExtensionsList(hwInfo);
+        OpenClCFeaturesContainer openclCFeatures;
+        if (requiresOpenClCFeatures(options)) {
+            getOpenclCFeaturesList(hwInfo, openclCFeatures);
+        }
+        auto compilerExtensions = convertEnabledExtensionsToCompilerInternalOptions(extensionsList.c_str(), openclCFeatures);
+        CompilerOptions::concatenateAppend(internalOptions, compilerExtensions);
+    }
+
     parseDebugSettings();
 
     // set up the device inside the program
@@ -585,28 +605,6 @@ int OfflineCompiler::parseCommandLine(size_t numArgs, const std::vector<std::str
         } else if (!argHelper->fileExists(inputFile)) {
             argHelper->printf("Error: Input file %s missing.\n", inputFile.c_str());
             retVal = INVALID_FILE;
-        } else {
-            retVal = deviceName.empty() ? 0 : getHardwareInfo(deviceName.c_str());
-            if (retVal != 0) {
-                argHelper->printf("Error: Cannot get HW Info for device %s.\n", deviceName.c_str());
-            } else if (false == deviceName.empty()) {
-                auto oclVersion = getOclVersionCompilerInternalOption(hwInfo.capabilityTable.clVersionSupport);
-                CompilerOptions::concatenateAppend(internalOptions, oclVersion);
-
-                std::string extensionsList = getExtensionsList(hwInfo);
-                if (requiresOpenClCFeatures(options)) {
-                    OpenClCFeaturesContainer openclCFeatures;
-                    getOpenclCFeaturesList(hwInfo, openclCFeatures);
-                    auto compilerExtensions = convertEnabledExtensionsToCompilerInternalOptions(extensionsList.c_str(), openclCFeatures);
-                    CompilerOptions::concatenateAppend(internalOptions, compilerExtensions);
-                } else {
-                    OpenClCFeaturesContainer emptyOpenClCFeatures;
-                    auto compilerExtensions = convertEnabledExtensionsToCompilerInternalOptions(extensionsList.c_str(), emptyOpenClCFeatures);
-                    CompilerOptions::concatenateAppend(internalOptions, compilerExtensions);
-                }
-            } else {
-                this->internalOptions = CompilerOptions::concatenate("-cl-ext=-all,+cl_khr_3d_image_writes", this->internalOptions);
-            }
         }
     }
 

shared/source/built_ins/builtins_binary.cmake

@@ -19,28 +19,28 @@ set(GENERATED_BUILTINS
     "aux_translation"
     "copy_buffer_rect"
     "copy_buffer_to_buffer"
-    "copy_buffer_to_image3d"
     "copy_kernel_timestamps"
     "fill_buffer"
+)
+
+set(GENERATED_BUILTINS_IMAGES
+    "copy_buffer_to_image3d"
+    "copy_image3d_to_buffer"
+    "copy_image_to_image1d"
+    "copy_image_to_image2d"
+    "copy_image_to_image3d"
     "fill_image1d"
     "fill_image2d"
     "fill_image3d"
 )
 
-set(GENERATED_BUILTINS_IMAGES
-    "copy_image3d_to_buffer"
-    "copy_image_to_image1d"
-    "copy_image_to_image2d"
-    "copy_image_to_image3d"
-)
-
 set(GENERATED_BUILTINS_IMAGES_STATELESS
+    "copy_buffer_to_image3d_stateless"
     "copy_image3d_to_buffer_stateless"
 )
 
 set(GENERATED_BUILTINS_STATELESS
     "copy_buffer_to_buffer_stateless"
-    "copy_buffer_to_image3d_stateless"
     "copy_buffer_rect_stateless"
     "fill_buffer_stateless"
 )

shared/test/unit_test/CMakeLists.txt

@@ -53,12 +53,14 @@ if(NOT SKIP_NEO_UNIT_TESTS AND NOT SKIP_UNIT_TESTS)
                  ${NEO_SOURCE_DIR}/opencl/test/unit_test/libult/create_command_stream.cpp
                  ${NEO_SOURCE_DIR}/opencl/test/unit_test/libult/io_functions.cpp
                  ${NEO_SOURCE_DIR}/opencl/test/unit_test/mocks/mock_platform.cpp
-                 ${NEO_SOURCE_DIR}/opencl/test/unit_test/ult_configuration.cpp
+                 ${NEO_SOURCE_DIR}/opencl/test/unit_test/test_macros/test_checks_ocl.cpp
                  ${NEO_SOURCE_DIR}/opencl/test/unit_test/ult_config_listener.cpp
+                 ${NEO_SOURCE_DIR}/opencl/test/unit_test/ult_configuration.cpp
                  ${NEO_SOURCE_DIR}/shared/source/helpers/allow_deferred_deleter.cpp
                  ${NEO_SOURCE_DIR}/shared/test/unit_test/helpers/api_specific_config_shared_tests.cpp
-                 ${NEO_SOURCE_DIR}/shared/test/unit_test/helpers/memory_management.cpp
                  ${NEO_SOURCE_DIR}/shared/test/unit_test/helpers/memory_leak_listener.cpp
+                 ${NEO_SOURCE_DIR}/shared/test/unit_test/helpers/memory_management.cpp
+                 ${NEO_SOURCE_DIR}/shared/test/unit_test/test_macros/test_checks_shared.cpp
                  $<TARGET_OBJECTS:mock_gmm>
   )
 

shared/test/unit_test/main.cpp

@@ -15,6 +15,7 @@
 #include "shared/test/unit_test/helpers/memory_leak_listener.h"
 #include "shared/test/unit_test/helpers/test_files.h"
 #include "shared/test/unit_test/helpers/ult_hw_config.inl"
+#include "shared/test/unit_test/test_macros/test_checks_shared.h"
 #include "shared/test/unit_test/tests_configuration.h"
 
 #include "opencl/source/os_interface/ocl_reg_path.h"
@@ -380,8 +381,14 @@ int main(int argc, char **argv) {
     MockCompilerDebugVars fclDebugVars;
     MockCompilerDebugVars igcDebugVars;
 
-    retrieveBinaryKernelFilename(fclDebugVars.fileName, KernelBinaryHelper::BUILT_INS + "_", ".bc");
-    retrieveBinaryKernelFilename(igcDebugVars.fileName, KernelBinaryHelper::BUILT_INS + "_", ".gen");
+    std::string builtInsFileName;
+    if (TestChecks::supportsImages(defaultHwInfo)) {
+        builtInsFileName = KernelBinaryHelper::BUILT_INS_WITH_IMAGES;
+    } else {
+        builtInsFileName = KernelBinaryHelper::BUILT_INS;
+    }
+    retrieveBinaryKernelFilename(fclDebugVars.fileName, builtInsFileName + "_", ".bc");
+    retrieveBinaryKernelFilename(igcDebugVars.fileName, builtInsFileName + "_", ".gen");
 
     gEnvironment->setMockFileNames(fclDebugVars.fileName, igcDebugVars.fileName);
     gEnvironment->setDefaultDebugVars(fclDebugVars, igcDebugVars, hwInfoForTests);

shared/test/unit_test/test_files/13936718464051189346.images.cl

@@ -158,6 +158,90 @@ __kernel void FillBufferSSHOffset(
     pDst[dstIndex] = pSrc[srcIndex];
 }
 
+//////////////////////////////////////////////////////////////////////////////
+__kernel void CopyBufferRectBytes2d(
+    __global const char* src,
+    __global char* dst,
+    uint4 SrcOrigin,
+    uint4 DstOrigin,
+    uint2 SrcPitch,
+    uint2 DstPitch )
+
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
+    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
+
+    *( dst + LDstOffset )  = *( src + LSrcOffset ); 
+
+}
+//////////////////////////////////////////////////////////////////////////////
+__kernel void CopyBufferRectBytes3d(
+    __global const char* src, 
+    __global char* dst, 
+    uint4 SrcOrigin, 
+    uint4 DstOrigin, 
+    uint2 SrcPitch, 
+    uint2 DstPitch ) 
+ 
+{ 
+    int x = get_global_id(0); 
+    int y = get_global_id(1); 
+    int z = get_global_id(2); 
+ 
+    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x ) + ( ( z + SrcOrigin.z ) * SrcPitch.y ); 
+    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x ) + ( ( z + DstOrigin.z ) * DstPitch.y ); 
+ 
+    *( dst + LDstOffset )  = *( src + LSrcOffset );  
+ 
+}
+
+__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
+    uint gid = get_global_id(0);
+    const ulong tsMask = (1ull << 32) - 1;
+    uint currentOffset = gid * 4;
+    dst[currentOffset] = 0;
+    dst[currentOffset + 1] = 0;
+    dst[currentOffset + 2] = 0;
+    dst[currentOffset + 3] = 0;
+
+    ulong srcPtr = srcEvents[gid];
+    __global uint *src = (__global uint *) srcPtr;
+    dst[currentOffset] = src[1] & tsMask;
+    dst[currentOffset + 1] = src[3] & tsMask;
+    if (useOnlyGlobalTimestamps != 0) {
+        dst[currentOffset + 2] = src[1] & tsMask;
+        dst[currentOffset + 3] = src[3] & tsMask;
+    } else {
+        dst[currentOffset + 2] = src[0] & tsMask;
+        dst[currentOffset + 3] = src[2] & tsMask;
+    }
+}
+
+__kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __global ulong* dst, __global ulong *offsets, uint useOnlyGlobalTimestamps) {
+    uint gid = get_global_id(0);
+    const ulong tsMask = (1ull << 32) - 1;
+    uint currentOffset = offsets[gid] / 8;
+    dst[currentOffset] = 0;
+    dst[currentOffset + 1] = 0;
+    dst[currentOffset + 2] = 0;
+    dst[currentOffset + 3] = 0;
+
+    ulong srcPtr = srcEvents[gid];
+    __global uint *src = (__global uint *) srcPtr;
+    dst[currentOffset] = src[1] & tsMask;
+    dst[currentOffset + 1] = src[3] & tsMask;
+    if (useOnlyGlobalTimestamps != 0) {
+        dst[currentOffset + 2] = src[1] & tsMask;
+        dst[currentOffset + 3] = src[3] & tsMask;
+    } else {
+        dst[currentOffset + 2] = src[0] & tsMask;
+        dst[currentOffset + 3] = src[2] & tsMask;
+    }
+}
+
 __kernel void FillImage1d(
     __write_only image1d_t output,
     uint4 color,
@@ -237,46 +321,6 @@ __kernel void CopyImageToImage3d(
     write_imageui(output, dstCoord, c);
 }
 
-//////////////////////////////////////////////////////////////////////////////
-__kernel void CopyBufferRectBytes2d(
-    __global const char* src,
-    __global char* dst,
-    uint4 SrcOrigin,
-    uint4 DstOrigin,
-    uint2 SrcPitch,
-    uint2 DstPitch )
-
-{
-    int x = get_global_id(0);
-    int y = get_global_id(1);
-
-    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
-    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
-
-    *( dst + LDstOffset )  = *( src + LSrcOffset ); 
-
-}
-//////////////////////////////////////////////////////////////////////////////
-__kernel void CopyBufferRectBytes3d(
-    __global const char* src, 
-    __global char* dst, 
-    uint4 SrcOrigin, 
-    uint4 DstOrigin, 
-    uint2 SrcPitch, 
-    uint2 DstPitch ) 
- 
-{ 
-    int x = get_global_id(0); 
-    int y = get_global_id(1); 
-    int z = get_global_id(2); 
- 
-    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x ) + ( ( z + SrcOrigin.z ) * SrcPitch.y ); 
-    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x ) + ( ( z + DstOrigin.z ) * DstPitch.y ); 
- 
-    *( dst + LDstOffset )  = *( src + LSrcOffset );  
- 
-}
-
 #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
 
 __kernel void CopyBufferToImage3dBytes(__global uchar *src,
@@ -560,47 +604,3 @@ __kernel void CopyImage3dToBuffer16Bytes(__read_only image3d_t input,
         *(__global uint4*)(dst + DstOffset + x * 16) = c;
     }
 }
-
-__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
-    uint gid = get_global_id(0);
-    const ulong tsMask = (1ull << 32) - 1;
-    uint currentOffset = gid * 4;
-    dst[currentOffset] = 0;
-    dst[currentOffset + 1] = 0;
-    dst[currentOffset + 2] = 0;
-    dst[currentOffset + 3] = 0;
-
-    ulong srcPtr = srcEvents[gid];
-    __global uint *src = (__global uint *) srcPtr;
-    dst[currentOffset] = src[1] & tsMask;
-    dst[currentOffset + 1] = src[3] & tsMask;
-    if (useOnlyGlobalTimestamps != 0) {
-        dst[currentOffset + 2] = src[1] & tsMask;
-        dst[currentOffset + 3] = src[3] & tsMask;
-    } else {
-        dst[currentOffset + 2] = src[0] & tsMask;
-        dst[currentOffset + 3] = src[2] & tsMask;
-    }
-}
-
-__kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __global ulong* dst, __global ulong *offsets, uint useOnlyGlobalTimestamps) {
-    uint gid = get_global_id(0);
-    const ulong tsMask = (1ull << 32) - 1;
-    uint currentOffset = offsets[gid] / 8;
-    dst[currentOffset] = 0;
-    dst[currentOffset + 1] = 0;
-    dst[currentOffset + 2] = 0;
-    dst[currentOffset + 3] = 0;
-
-    ulong srcPtr = srcEvents[gid];
-    __global uint *src = (__global uint *) srcPtr;
-    dst[currentOffset] = src[1] & tsMask;
-    dst[currentOffset + 1] = src[3] & tsMask;
-    if (useOnlyGlobalTimestamps != 0) {
-        dst[currentOffset + 2] = src[1] & tsMask;
-        dst[currentOffset + 3] = src[3] & tsMask;
-    } else {
-        dst[currentOffset + 2] = src[0] & tsMask;
-        dst[currentOffset + 3] = src[2] & tsMask;
-    }
-}

shared/test/unit_test/test_files/17475704436818206941.cl

@@ -0,0 +1,243 @@
+/*
+ * Copyright (C) 2020 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+__kernel void fullCopy(__global const uint* src, __global uint* dst) {
+    unsigned int gid = get_global_id(0);
+    uint4 loaded = vload4(gid, src);
+    vstore4(loaded, gid, dst);
+}
+
+__kernel void CopyBufferToBufferBytes(
+    const __global uchar* pSrc,
+    __global uchar* pDst,
+    uint srcOffsetInBytes,
+    uint dstOffsetInBytes,
+    uint bytesToRead )
+{
+    pSrc += ( srcOffsetInBytes + get_global_id(0) );
+    pDst += ( dstOffsetInBytes + get_global_id(0) );
+    pDst[ 0 ] = pSrc[ 0 ];
+}
+
+__kernel void CopyBufferToBufferLeftLeftover(
+    const __global uchar* pSrc,
+    __global uchar* pDst,
+    uint srcOffsetInBytes,
+    uint dstOffsetInBytes)
+{
+    unsigned int gid = get_global_id(0);
+    pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ];
+}
+
+__kernel void CopyBufferToBufferMiddle(
+    const __global uint* pSrc,
+    __global uint* pDst,
+    uint srcOffsetInBytes,
+    uint dstOffsetInBytes)
+{
+    unsigned int gid = get_global_id(0);
+    pDst += dstOffsetInBytes >> 2;
+    pSrc += srcOffsetInBytes >> 2;
+    uint4 loaded = vload4(gid, pSrc);
+    vstore4(loaded, gid, pDst);
+}
+
+__kernel void CopyBufferToBufferRightLeftover(
+    const __global uchar* pSrc,
+    __global uchar* pDst,
+    uint srcOffsetInBytes,
+    uint dstOffsetInBytes)
+{
+    unsigned int gid = get_global_id(0);
+    pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ];
+}
+
+__kernel void copyBufferToBufferBytesSingle(__global uchar *dst, const __global uchar *src) {
+    unsigned int gid = get_global_id(0);
+    dst[gid] = (uchar)(src[gid]);
+}
+__kernel void CopyBufferToBufferSideRegion(
+    __global uchar* pDst,
+    const __global uchar* pSrc,
+    unsigned int len,
+    uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
+    uint srcSshOffset // Offset needed in case ptr has been adjusted for SSH alignment
+    )
+{
+    unsigned int gid = get_global_id(0);
+    __global uchar* pDstWithOffset = (__global uchar*)((__global uchar*)pDst + dstSshOffset);
+    __global uchar* pSrcWithOffset = (__global uchar*)((__global uchar*)pSrc + srcSshOffset);
+    if (gid < len) {
+        pDstWithOffset[ gid ] = pSrcWithOffset[ gid ];
+    }
+}
+
+__kernel void CopyBufferToBufferMiddleRegion(
+    __global uint* pDst,
+    const __global uint* pSrc,
+    unsigned int elems,
+    uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
+    uint srcSshOffset // Offset needed in case ptr has been adjusted for SSH alignment
+    )
+{
+    unsigned int gid = get_global_id(0);
+    __global uint* pDstWithOffset = (__global uint*)((__global uchar*)pDst + dstSshOffset);
+    __global uint* pSrcWithOffset = (__global uint*)((__global uchar*)pSrc + srcSshOffset);
+    if (gid < elems) {
+        uint4 loaded = vload4(gid, pSrcWithOffset);
+        vstore4(loaded, gid, pDstWithOffset);
+    }
+}
+
+// assumption is local work size = pattern size
+__kernel void FillBufferBytes(
+    __global uchar* pDst,
+    uint dstOffsetInBytes,
+    const __global uchar* pPattern )
+{
+    uint dstIndex = get_global_id(0) + dstOffsetInBytes;
+    uint srcIndex = get_local_id(0);
+    pDst[dstIndex] = pPattern[srcIndex];
+}
+
+__kernel void FillBufferLeftLeftover(
+    __global uchar* pDst,
+    uint dstOffsetInBytes,
+    const __global uchar* pPattern,
+    const uint patternSizeInEls )
+{
+    uint gid = get_global_id(0);
+    pDst[ gid + dstOffsetInBytes ] = pPattern[ gid & (patternSizeInEls - 1) ];
+}
+
+__kernel void FillBufferMiddle(
+    __global uchar* pDst,
+    uint dstOffsetInBytes,
+    const __global uint* pPattern,
+    const uint patternSizeInEls )
+{
+    uint gid = get_global_id(0);
+    ((__global uint*)(pDst + dstOffsetInBytes))[gid] = pPattern[ gid & (patternSizeInEls - 1) ];
+}
+
+__kernel void FillBufferRightLeftover(
+    __global uchar* pDst,
+    uint dstOffsetInBytes,
+    const __global uchar* pPattern,
+    const uint patternSizeInEls )
+{
+    uint gid = get_global_id(0);
+    pDst[ gid + dstOffsetInBytes ] = pPattern[ gid & (patternSizeInEls - 1) ];
+}
+
+__kernel void FillBufferImmediate(
+    __global uchar* ptr,
+    uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
+    const uint value)
+{
+    uint dstIndex = get_global_id(0);
+    __global uchar* pDst = (__global uchar*)ptr + dstSshOffset;
+    pDst[dstIndex] = value;
+}
+
+__kernel void FillBufferSSHOffset(
+    __global uchar* ptr,
+    uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
+    const __global uchar* pPattern,
+    uint patternSshOffset // Offset needed in case pPattern has been adjusted for SSH alignment
+)
+{
+    uint dstIndex = get_global_id(0);
+    uint srcIndex = get_local_id(0);
+    __global uchar* pDst = (__global uchar*)ptr + dstSshOffset;
+    __global uchar* pSrc = (__global uchar*)pPattern + patternSshOffset;
+    pDst[dstIndex] = pSrc[srcIndex];
+}
+
+//////////////////////////////////////////////////////////////////////////////
+__kernel void CopyBufferRectBytes2d(
+    __global const char* src,
+    __global char* dst,
+    uint4 SrcOrigin,
+    uint4 DstOrigin,
+    uint2 SrcPitch,
+    uint2 DstPitch )
+
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
+    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
+
+    *( dst + LDstOffset )  = *( src + LSrcOffset ); 
+
+}
+//////////////////////////////////////////////////////////////////////////////
+__kernel void CopyBufferRectBytes3d(
+    __global const char* src, 
+    __global char* dst, 
+    uint4 SrcOrigin, 
+    uint4 DstOrigin, 
+    uint2 SrcPitch, 
+    uint2 DstPitch ) 
+ 
+{ 
+    int x = get_global_id(0); 
+    int y = get_global_id(1); 
+    int z = get_global_id(2); 
+ 
+    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x ) + ( ( z + SrcOrigin.z ) * SrcPitch.y ); 
+    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x ) + ( ( z + DstOrigin.z ) * DstPitch.y ); 
+ 
+    *( dst + LDstOffset )  = *( src + LSrcOffset );  
+ 
+}
+
+__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
+    uint gid = get_global_id(0);
+    const ulong tsMask = (1ull << 32) - 1;
+    uint currentOffset = gid * 4;
+    dst[currentOffset] = 0;
+    dst[currentOffset + 1] = 0;
+    dst[currentOffset + 2] = 0;
+    dst[currentOffset + 3] = 0;
+
+    ulong srcPtr = srcEvents[gid];
+    __global uint *src = (__global uint *) srcPtr;
+    dst[currentOffset] = src[1] & tsMask;
+    dst[currentOffset + 1] = src[3] & tsMask;
+    if (useOnlyGlobalTimestamps != 0) {
+        dst[currentOffset + 2] = src[1] & tsMask;
+        dst[currentOffset + 3] = src[3] & tsMask;
+    } else {
+        dst[currentOffset + 2] = src[0] & tsMask;
+        dst[currentOffset + 3] = src[2] & tsMask;
+    }
+}
+
+__kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __global ulong* dst, __global ulong *offsets, uint useOnlyGlobalTimestamps) {
+    uint gid = get_global_id(0);
+    const ulong tsMask = (1ull << 32) - 1;
+    uint currentOffset = offsets[gid] / 8;
+    dst[currentOffset] = 0;
+    dst[currentOffset + 1] = 0;
+    dst[currentOffset + 2] = 0;
+    dst[currentOffset + 3] = 0;
+
+    ulong srcPtr = srcEvents[gid];
+    __global uint *src = (__global uint *) srcPtr;
+    dst[currentOffset] = src[1] & tsMask;
+    dst[currentOffset + 1] = src[3] & tsMask;
+    if (useOnlyGlobalTimestamps != 0) {
+        dst[currentOffset + 2] = src[1] & tsMask;
+        dst[currentOffset + 3] = src[3] & tsMask;
+    } else {
+        dst[currentOffset + 2] = src[0] & tsMask;
+        dst[currentOffset + 3] = src[2] & tsMask;
+    }
+}

shared/test/unit_test/test_files/17475704436818206941_options.txt

@@ -0,0 +1,8 @@
+/*
+ * Copyright (C) 2020 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+-cl-fast-relaxed-math

shared/test/unit_test/test_files/CMakeLists.txt

@@ -67,6 +67,7 @@ function(neo_shared_copy_test_files target product revision_id)
 endfunction()
 
 file(GLOB_RECURSE TEST_KERNELS *.cl)
+file(GLOB_RECURSE TEST_KERNELS_IMAGES *.images.cl)
 add_custom_target(prepare_test_kernel_for_shared)
 set_target_properties(prepare_test_kernel_for_shared PROPERTIES FOLDER "${SHARED_TEST_PROJECTS_FOLDER}/${SHARED_TEST_PROJECTS_SUB_FOLDER}")
 
@@ -87,10 +88,17 @@ macro(macro_for_each_gen)
 
     if(${GEN_TYPE}_HAS_${PLATFORM_TYPE})
 
+      set(KERNELS_TO_COMPILE ${TEST_KERNELS})
+      set(IMAGE_SUPPORT FALSE)
+      GEN_CONTAINS_PLATFORMS("SUPPORTED_IMAGES" ${GEN_TYPE} IMAGE_SUPPORT)
+      if(NOT IMAGE_SUPPORT)
+        list(REMOVE_ITEM KERNELS_TO_COMPILE ${TEST_KERNELS_IMAGES})
+      endif()
+
       get_family_name_with_type(${GEN_TYPE} ${PLATFORM_TYPE})
       set(PLATFORM_LOWER ${DEFAULT_SUPPORTED_${GEN_TYPE}_${PLATFORM_TYPE}_PLATFORM})
       foreach(REVISION_ID ${${PLATFORM_TYPE}_${GEN_TYPE}_REVISIONS})
-        compile_kernels_gen(${family_name_with_type} ${REVISION_ID} ${PLATFORM_LOWER} ${family_name_with_type} ${TEST_KERNELS})
+        compile_kernels_gen(${family_name_with_type} ${REVISION_ID} ${PLATFORM_LOWER} ${family_name_with_type} ${KERNELS_TO_COMPILE})
 
         add_custom_target(prepare_test_kernel_for_shared_${family_name_with_type}_${REVISION_ID} DEPENDS ${compiled_kernels_${family_name_with_type}_${REVISION_ID}})
         set_target_properties(prepare_test_kernel_for_shared_${family_name_with_type}_${REVISION_ID} PROPERTIES FOLDER "${SHARED_TEST_PROJECTS_FOLDER}/${SHARED_TEST_PROJECTS_SUB_FOLDER}/${PLATFORM_SPECIFIC_TEST_TARGETS_FOLDER}/${family_name_with_type}/${REVISION_ID}")

shared/test/unit_test/test_macros/test_checks_shared.cpp

@@ -8,6 +8,7 @@
 #include "shared/test/unit_test/test_macros/test_checks_shared.h"
 
 #include "shared/source/device/device.h"
+#include "shared/source/helpers/constants.h"
 #include "shared/source/helpers/hw_helper.h"
 #include "shared/test/unit_test/helpers/default_hw_info.h"
 
@@ -15,7 +16,11 @@
 
 using namespace NEO;
 
-bool NEO::TestChecks::supportsBlitter(const HardwareInfo *pHardwareInfo) {
+bool TestChecks::is64Bit() {
+    return ::is64bit;
+}
+
+bool TestChecks::supportsBlitter(const HardwareInfo *pHardwareInfo) {
     auto engines = HwHelper::get(::renderCoreFamily).getGpgpuEngineInstances(*pHardwareInfo);
     for (const auto &engine : engines) {
         if (engine.first == aub_stream::EngineType::ENGINE_BCS) {
@@ -25,6 +30,14 @@ bool NEO::TestChecks::supportsBlitter(const HardwareInfo *pHardwareInfo) {
     return false;
 }
 
+bool TestChecks::supportsImages(const HardwareInfo &hardwareInfo) {
+    return hardwareInfo.capabilityTable.supportsImages;
+}
+
+bool TestChecks::supportsImages(const std::unique_ptr<HardwareInfo> &pHardwareInfo) {
+    return supportsImages(*pHardwareInfo);
+}
+
 bool TestChecks::supportsSvm(const HardwareInfo *pHardwareInfo) {
     return pHardwareInfo->capabilityTable.ftrSvm;
 }
@@ -38,7 +51,7 @@ bool TestChecks::supportsSvm(const Device *pDevice) {
 class TestMacrosIfNotMatchTearDownCall : public ::testing::Test {
   public:
     void expectCorrectPlatform() {
-        EXPECT_EQ(IGFX_SKYLAKE, NEO::defaultHwInfo->platform.eProductFamily);
+        EXPECT_EQ(IGFX_SKYLAKE, defaultHwInfo->platform.eProductFamily);
     }
     void SetUp() override {
         expectCorrectPlatform();

shared/test/unit_test/test_macros/test_checks_shared.h

@@ -15,7 +15,10 @@ class Device;
 struct HardwareInfo;
 
 namespace TestChecks {
+bool is64Bit();
 bool supportsBlitter(const HardwareInfo *pHardwareInfo);
+bool supportsImages(const HardwareInfo &hardwareInfo);
+bool supportsImages(const std::unique_ptr<HardwareInfo> &pHardwareInfo);
 bool supportsSvm(const HardwareInfo *pHardwareInfo);
 bool supportsSvm(const std::unique_ptr<HardwareInfo> &pHardwareInfo);
 bool supportsSvm(const Device *pDevice);
@@ -23,6 +26,11 @@ bool supportsSvm(const Device *pDevice);
 
 } // namespace NEO
 
+#define REQUIRE_64BIT_OR_SKIP()                \
+    if (NEO::TestChecks::is64Bit() == false) { \
+        GTEST_SKIP();                          \
+    }
+
 #define REQUIRE_SVM_OR_SKIP(param)                      \
     if (NEO::TestChecks::supportsSvm(param) == false) { \
         GTEST_SKIP();                                   \
@@ -32,3 +40,8 @@ bool supportsSvm(const Device *pDevice);
     if (NEO::TestChecks::supportsBlitter(param) == false) { \
         GTEST_SKIP();                                       \
     }
+
+#define REQUIRE_IMAGES_OR_SKIP(param)                      \
+    if (NEO::TestChecks::supportsImages(param) == false) { \
+        GTEST_SKIP();                                      \
+    }