[SYCL][ESIMD] Don't add RegisterAllocMode metadata for ESIMD kernels

This works around a VC issue.

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

Author: sarnex

llvm/lib/SYCLLowerIR/LowerKernelProps.cpp

@@ -50,7 +50,8 @@ void processSetKernelPropertiesCall(CallInst &CI) {
       GraphNode->addFnAttr(::sycl::kernel_props::ATTR_LARGE_GRF);
       // Add RegisterAllocMode metadata with arg 2 to the kernel to tell
       // IGC to compile this kernel in large GRF mode. 2 means large.
-      if (GraphNode->getCallingConv() == CallingConv::SPIR_KERNEL) {
+      if (GraphNode->getCallingConv() == CallingConv::SPIR_KERNEL &&
+          !GraphNode->hasMetadata("sycl_explicit_simd")) {
         auto &Ctx = GraphNode->getContext();
         Metadata *AttrMDArgs[] = {ConstantAsMetadata::get(
             Constant::getIntegerValue(Type::getInt32Ty(Ctx), APInt(32, 2)))};

llvm/test/SYCLLowerIR/lower_kernel_props.ll

@@ -5,7 +5,7 @@
 ; - remove the intrinsic call 
 ; - mark the kernel with corresponding attribute (only "large-grf" for now)
 
-; RUN: opt -passes=lower-kernel-props -S %s -o - | FileCheck %s
+; RUN: opt -passes=lower-kernel-props -S %s -o - | FileCheck %s --implicit-check-not='RegisterAllocMode'
 
 ; ModuleID = 'large_grf.bc'
 source_filename = "llvm-link"
@@ -33,8 +33,7 @@ define weak_odr dso_local spir_kernel void @__large_grf_kernel1() !sycl_explicit
 
 ; -- This kernel calls the marker function directly
 define weak_odr dso_local spir_kernel void @__large_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
-; CHECK: {{.*}} spir_kernel void @__large_grf_kernel2() #0 {{.*}} !RegisterAllocMode ![[MetadataArg:[0-9]+]]
-; CHECK: ![[MetadataArg]] = !{i32 2}
+; CHECK: {{.*}} spir_kernel void @__large_grf_kernel2() #0 {{.*}}
   call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
   ret void
 }

llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll

@@ -9,7 +9,7 @@
 
 ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table
 ; RUN: FileCheck %s -input-file=%t.table
-; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
+; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR  --implicit-check-not='RegisterAllocMode'
 ; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
 ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM
 ; RUN: FileCheck %s -input-file=%t_esimd_1.sym --check-prefixes CHECK-ESIMD-SYM
@@ -58,8 +58,7 @@ entry:
 declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef)
 
 define weak_odr dso_local spir_kernel void @__ESIMD_large_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
-; CHECK-ESIMD-LargeGRF-IR: @__ESIMD_large_grf_kernel() {{.*}} !RegisterAllocMode ![[MetadataArg:[0-9]+]]
-; CHECK-ESIMD-LargeGRF-IR: ![[MetadataArg]] = !{i32 2}
+; CHECK-ESIMD-LargeGRF-IR: @__ESIMD_large_grf_kernel() {{.*}}
 entry:
   call spir_func void @_Z17large_grf_markerv()
   ret void