enable RemoveUnusedIdImplicitArguments for Xe family

RemoveUnusedIdImplicitArguments is an option to remove global_id_offset
and enqueued_local_size from kernel's implicit argument if they
are unused (true if kernel does not use get_global_id()).

Enable this option by default for Xe family.

Author: pkwasnie-intel

IGC/Compiler/CISACodeGen/OpenCLKernelCodeGen.cpp

@@ -2454,7 +2454,7 @@ namespace IGC
                 arg.getArgType() == KernelArg::ArgType::IMPLICIT_BUFFER_SIZE) &&
                 arg.getArg()->use_empty();
 
-            if (IGC_IS_FLAG_ENABLED(RemoveUnusedIdImplicitArguments))
+            if (m_Context->platform.allowRemovingUnusedImplicitArguments())
             {
                 IsUnusedArg |=
                     (arg.getArgType() == KernelArg::ArgType::IMPLICIT_PAYLOAD_HEADER || // contains global_id_offset

IGC/Compiler/CISACodeGen/Platform.hpp

@@ -1897,5 +1897,13 @@ bool allowProceedBasedApproachForRayQueryDynamicRayManagementMechanism() const
     return IGC_IS_FLAG_DISABLED(DisableProceedBasedApproachForRayQueryDynamicRayManagementMechanism);
 }
 
+bool allowRemovingUnusedImplicitArguments() const
+{
+    if (IGC_IS_FLAG_SET(RemoveUnusedIdImplicitArguments))
+        return IGC_IS_FLAG_ENABLED(RemoveUnusedIdImplicitArguments);
+
+    return isCoreChildOf(IGFX_XE_HP_CORE) && !isCoreChildOf(IGFX_XE2_HPG_CORE);
+}
+
 };
 }//namespace IGC