[ESIMD] Add slm_init<compile_time_constant>() version. (#6083)

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Author: kbobrovs

sycl/include/sycl/ext/intel/esimd/memory.hpp

@@ -798,6 +798,13 @@ __ESIMD_API void barrier() {
 /// @{
 
 /// Declare per-work-group slm size.
+/// @tparam SLMSize  Shared Local Memory (SLM) size
+template <uint32_t SLMSize> __ESIMD_API void slm_init() {
+  __esimd_slm_init(SLMSize);
+}
+
+/// Declare per-work-group slm size. Non-constant argument version to be used
+/// with specialization constants only.
 /// @param size  Shared Local Memory (SLM) size
 __ESIMD_API void slm_init(uint32_t size) { __esimd_slm_init(size); }
 

sycl/test/esimd/genx_func_attr.cpp

@@ -3,7 +3,6 @@
 // RUN: FileCheck %s -input-file=%t_esimd_0.ll
 
 // Checks ESIMD intrinsic translation.
-// NOTE: must be run in -O0, as optimizer optimizes away some of the code
 
 #include <CL/sycl.hpp>
 #include <CL/sycl/detail/image_ocl_types.hpp>
@@ -17,7 +16,7 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
 }
 
 SYCL_ESIMD_FUNCTION SYCL_EXTERNAL ESIMD_NOINLINE void callee(int x) {
-  slm_init(1234);
+  slm_init<1234>();
   sycl::ext::intel::experimental::esimd::named_barrier_init<13>();
 }
 
@@ -30,7 +29,7 @@ void caller_abc(int x) {
 // inherits only NBarrierCount from callee
 void caller_xyz(int x) {
   kernel<class kernel_xyz>([=]() SYCL_ESIMD_KERNEL {
-    slm_init(1235);
+    slm_init(1235); // also works in non-O0
     callee(x);
   });
   // CHECK: define dso_local spir_kernel void @_ZTSZ10caller_xyziE10kernel_xyz(i32 noundef "VCArgumentIOKind"="0" %{{.*}}) local_unnamed_addr #3

sycl/test/esimd/slm_gather_scatter_rgba.cpp

@@ -11,7 +11,7 @@ void caller() SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, sizeof(int) * 4);
   simd<int, 128> v1(0, 1);
 
-  slm_init(1024);
+  slm_init<1024>();
 
   auto v0 = slm_gather_rgba<int, 32, rgba_channel_mask::ABGR>(offsets);