Lower memcpy right after LowerByValAttribute pass

`LowerByValAttribute` inserts calls to memcpy intrinsic. This change introduces
calling `ReplaceUnsupportedIntrinsics` pass right after `LowerByValAttribute`
to transform all `memcpy` calls to a series a loads/stores. Without that,
`LegalizationPass` doesn't legalize all loads and stores, as they are in a form
of `memcpy` calls.

This change also implements the possibility to run `LowerByValAttribute` pass
only if optimizations are enabled. `LowerByValAttribute` pass generates an
explicit copy (alloca+memcpy) for all function arguments with `byval` attribute.
The copy is not necessary if a SPIRV was compiled in O0, as in such case,
FE compilers will generate an explicit copy (not optimize it out) in SPIRV.

Author: aratajew

IGC/Compiler/CISACodeGen/ShaderCodeGen.cpp

@@ -295,8 +295,13 @@ void AddAnalysisPasses(CodeGenContext& ctx, IGCPassManager& mpm)
             }
         }
         mpm.add(createPromoteMemoryToRegisterPass());
-        if (ctx.type == ShaderType::OPENCL_SHADER)
+        if (ctx.type == ShaderType::OPENCL_SHADER &&
+            !isOptDisabled &&
+            IGC_IS_FLAG_ENABLED(EnableExplicitCopyForByVal))
+        {
             mpm.add(new LowerByValAttribute());
+            mpm.add(createReplaceUnsupportedIntrinsicsPass());
+        }
         // Resolving private memory allocas
         mpm.add(CreatePrivateMemoryResolution());
     }
@@ -673,8 +678,13 @@ void AddLegalizationPasses(CodeGenContext& ctx, IGCPassManager& mpm, PSSignature
     if (!(IGC_IS_FLAG_ENABLED(EnableUnmaskedFunctions) &&
           IGC_IS_FLAG_ENABLED(LateInlineUnmaskedFunc)))
     {
-        if (ctx.type == ShaderType::OPENCL_SHADER)
+        if (ctx.type == ShaderType::OPENCL_SHADER &&
+            !isOptDisabled &&
+            IGC_IS_FLAG_ENABLED(EnableExplicitCopyForByVal))
+        {
             mpm.add(new LowerByValAttribute());
+            mpm.add(createReplaceUnsupportedIntrinsicsPass());
+        }
         mpm.add(CreatePrivateMemoryResolution());
     }
     // Should help MemOpt pass to merge more loads

IGC/Compiler/Optimizer/OpenCLPasses/PrivateMemory/LowerByValAttribute.cpp

@@ -8,7 +8,6 @@ SPDX-License-Identifier: MIT
 
 #include "Compiler/Optimizer/OpenCLPasses/PrivateMemory/LowerByValAttribute.hpp"
 #include "Compiler/IGCPassSupport.h"
-#include "common/igc_regkeys.hpp"
 
 #include "common/LLVMWarningsPush.hpp"
 #include <llvm/Pass.h>
@@ -49,9 +48,6 @@ LowerByValAttribute::LowerByValAttribute(void) : FunctionPass(ID)
 
 bool LowerByValAttribute::runOnFunction(Function& F)
 {
-    if (IGC_IS_FLAG_DISABLED(EnableExplicitCopyForByVal))
-        return false;
-
     visit(F);
 
     return m_changed;

IGC/Compiler/tests/LowerByValAttribute/explicit-copy-for-byval.ll

@@ -6,9 +6,7 @@
 ;
 ;============================ end_copyright_notice =============================
 ;
-; REQUIRES: regkeys
-;
-; RUN: igc_opt %s -S -o - --igc-lower-byval-attribute -regkey EnableExplicitCopyForByVal | FileCheck %s
+; RUN: igc_opt %s -S -o - --igc-lower-byval-attribute | FileCheck %s
 
 ; ------------------------------------------------
 ; LegalizeFunctionSignatures

IGC/ocloc_tests/lowerbyval-before-privatememresolution.cl

@@ -8,16 +8,18 @@ SPDX-License-Identifier: MIT
 
 // REQUIRES: regkeys
 
-// RUN: ocloc compile -file %s -options " -igc_opts 'ShaderDisplayAllPassesNames=1'"  -device dg2 2>&1 | FileCheck %s
+// RUN: ocloc compile -file %s -options " -igc_opts 'ShaderDisplayAllPassesNames=1 EnableExplicitCopyForByVal=1'"  -device dg2 2>&1 | FileCheck %s
 
 // Verify if LowerByValAttribute pass is run right before PrivateMemoryResolution pass.
+// (ReplaceUnsupportedIntrinsics is run between them just to lower memcpy instructions inserted by LowerByValAttribute)
 
 // If LowerByValAttribute was run after PrivateMemoryResolution, alloca instructions inserted by LowerByValAttribute
 // wouldn't be resolved.
 // If LowerByValAttribute was run earlier than right before PrivateMemoryResolution, there is a chance that
 // optimization passes would remove explicit copy (alloca + memcpy) inserted by LowerByValAttribute.
 
 // CHECK: LowerByValAttribute
+// CHECK-NEXT: ReplaceUnsupportedIntrinsics
 // CHECK-NEXT: PrivateMemoryResolution
 
 __kernel void foo(int a, int b, __global int *res) { *res = a + b; }