[SYCL][ESIMD] Fix an issue with setting VCSLMSize attribute (#9155)

With the recent SLM handling changes we started assigning
"VCSLMSize"="4294967295" attribute to the kernels using classic
slm_init() with the size taken from the spec constant. Previous behavior
is to set 0.

VC backend takes the size from the intrinsic if [the provided size is
greater than in
metadata.](https://github.com/intel/intel-graphics-compiler/blob/2aa8d536f8823833a18fc5f84c9b4bce9c8060dd/IGC/VectorCompiler/lib/GenXCodeGen/GenXLowering.cpp#L5662)
Kernel will not compile with max_int amount of requested SLM:

Author: turinevgeny

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp

@@ -635,6 +635,12 @@ size_t lowerSLMReservationCalls(Module &M) {
   // - now set each kernel's SLMSize metadata to the pre-calculated value
   for (auto &E : Kernel2MaxSLM) {
     int MaxSLM = E.second;
+    // Clamp negative values to 0. MaxSLM could have been not estimated, e.g.
+    // due to having __esimd_slm_init with non-const operand (specialization
+    // constant case). VC backend will use size provided in __esimd_slm_init
+    // if it is greater than value provided in metadata.
+    if (MaxSLM < 0)
+      MaxSLM = 0;
     llvm::Value *MaxSLMv =
         llvm::ConstantInt::get(Type::getInt32Ty(M.getContext()), MaxSLM);
     const Function *Kernel = E.first;

sycl/test/esimd/slm_init_specconst_size.cpp

@@ -0,0 +1,32 @@
+// RUN: %clangxx -O2 -fsycl -fsycl-device-only -Xclang -emit-llvm %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
+
+// Checks that we set 0 as VCSLMSize when slm_init is used with
+// non-constant operand, like with specialization constant.
+
+#include <sycl/detail/image_ocl_types.hpp>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl::ext::intel::esimd;
+using namespace sycl::ext::intel::experimental::esimd;
+
+constexpr sycl::specialization_id<uint64_t> Size(1024);
+
+int main() {
+  sycl::queue queue;
+  {
+    queue.submit([&](sycl::handler &cgh) {
+      cgh.single_task<class Kernel3Name>(
+          [=](sycl::kernel_handler kh) SYCL_ESIMD_KERNEL {
+            slm_init(kh.get_specialization_constant<Size>());
+          });
+      // CHECK: define weak_odr dso_local spir_kernel void @{{.*}}(i8 addrspace(1)* noundef align 1 "VCArgumentIOKind"="0" %{{.*}}) local_unnamed_addr #1
+    });
+  }
+
+  return 0;
+}
+
+// CHECK: attributes #1 = { {{.*}} "VCSLMSize"="0"