Remove InsertGenericPtrArithmeticMetadata pass

`InsertGenericPtrArithmeticMetadata` pass is responsible for adding
`generic.arith` metadata to addrspacecast from {`private`, `local`, `global`}
addrspace to `generic` addrspace. If it's present, then `EmitVISAPass` bypasses
tagging a generic pointer while emitting the addrspacecast. The metadata gets
added if `InsertGenericPtrArithmeticMetadata` pass detects that the addrspacecast's
result is used for pointer arithmetic.

Generic pointer tag can be problematic while comparing a generic pointer with
`NULL` or with some named pointer {`private`, `local`, `global`}. However,
in this case it's not necessary to depend on `generic.arith` metadata, because
IGC has a pattern match that detects generic pointers comparison and
resets tag bits if so.

I don't see any reasonable arguments behind keeping the mechanism for
`generic.arith` at the moment, therefore removing it. Precommit testing hasn't
detected any functional regression on this change.

Moreover, `generic.arith` may cause functional incorrectness. Here is the example:
```cpp
__kernel void generic_ptr_issue(local int* p0, local ulong* p1) {
  generic int* p0_g = p0;                      // addrspacecast local -> generic
  ulong p0_g_addr_plus_4 = (ulong)p0_g + 4;    // (not tagged due to generic.arith)
  p1[0] = p0_g_addr_plus_4;                    // untagged address written to local memory

  generic int* p0_g_plus_4 = (generic int*)p1[0];  // loading untagged address as generic
  p0_g_plus_4[0] = 5;                              // address is actually local, but global branch is taken due to a missing tag
};
```

Author: aratajew

IGC/Compiler/CISACodeGen/CMakeLists.txt

@@ -47,7 +47,6 @@ set(IGC_BUILD__SRC__CISACodeGen_Common
     "${CMAKE_CURRENT_SOURCE_DIR}/getCacheOpts.cpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/HalfPromotion.cpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/HoistCongruentPhi.cpp"
-    "${CMAKE_CURRENT_SOURCE_DIR}/InsertGenericPtrArithmeticMetadata.cpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/LdShrink.cpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/LiveVars.cpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/LivenessAnalysis.cpp"
@@ -149,7 +148,6 @@ set(IGC_BUILD__HDR__CISACodeGen_Common
     "${CMAKE_CURRENT_SOURCE_DIR}/getCacheOpts.h"
     "${CMAKE_CURRENT_SOURCE_DIR}/HalfPromotion.h"
     "${CMAKE_CURRENT_SOURCE_DIR}/HoistCongruentPhi.hpp"
-    "${CMAKE_CURRENT_SOURCE_DIR}/InsertGenericPtrArithmeticMetadata.hpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/LdShrink.h"
     "${CMAKE_CURRENT_SOURCE_DIR}/LiveVars.hpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/LivenessAnalysis.hpp"

IGC/Compiler/CISACodeGen/EmitVISAPass.cpp

@@ -10021,14 +10021,6 @@ void EmitPass::emitAddrSpaceCast(llvm::AddrSpaceCastInst* addrSpaceCast)
         // Address space cast is in the form of {private, local, global} -> generic
         // A tag is added according to the address space of the source
 
-        MDNode* genericMD = addrSpaceCast->getMetadata("generic.arith");
-        if (genericMD)
-        {
-            m_encoder->Cast(m_destination, srcV);
-            m_encoder->Push();
-            return;
-        }
-
         if (sourceAddrSpace == ADDRESS_SPACE_PRIVATE && (!m_pCtx->allocatePrivateAsGlobalBuffer() || m_pCtx->mustDistinguishBetweenPrivateAndGlobalPtr()))
         {
             emitAddrSpaceToGenericCast(addrSpaceCast, srcV, 1);

IGC/Compiler/CISACodeGen/InsertGenericPtrArithmeticMetadata.cpp

@@ -52,7 +52,6 @@ SPDX-License-Identifier: MIT
 #include "Compiler/CISACodeGen/UniformAssumptions.hpp"
 #include "Compiler/CISACodeGen/VectorProcess.hpp"
 #include "Compiler/CISACodeGen/RuntimeValueLegalizationPass.h"
-#include "Compiler/CISACodeGen/InsertGenericPtrArithmeticMetadata.hpp"
 #include "Compiler/CISACodeGen/LowerGEPForPrivMem.hpp"
 #include "Compiler/CISACodeGen/POSH_RemoveNonPositionOutput.h"
 #include "Compiler/CISACodeGen/RegisterEstimator.hpp"
@@ -1107,10 +1106,6 @@ void AddLegalizationPasses(CodeGenContext& ctx, IGCPassManager& mpm, PSSignature
 
     mpm.add(new WAFMinFMax());
 
-    // Preferred to be added after llvm instruction combining, otherwise 'generic.arith'
-    // metadata may get lost during optimizations.
-    mpm.add(new InsertGenericPtrArithmeticMetadata());
-
     mpm.add(createTimeStatsCounterPass(&ctx, TIME_CG_Legalization, STATS_COUNTER_END));
 
     COMPILER_TIME_END(&ctx, TIME_CG_Add_Legalization_Passes);

IGC/Compiler/CISACodeGen/InsertGenericPtrArithmeticMetadata.hpp

@@ -89,7 +89,6 @@ void initializeHandleLoadStoreInstructionsPass(llvm::PassRegistry&);
 void initializeIGCConstPropPass(llvm::PassRegistry&);
 void initializeJointMatrixFuncsResolutionPassPass(llvm::PassRegistry&);
 void initializeGatingSimilarSamplesPass(llvm::PassRegistry&);
-void initializeInsertGenericPtrArithmeticMetadataPass(llvm::PassRegistry&);
 void initializeImageFuncResolutionPass(llvm::PassRegistry&);
 void initializeImageFuncsAnalysisPass(llvm::PassRegistry&);
 void initializeImage3dToImage2darrayPass(llvm::PassRegistry&);