Address feedback

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -279,10 +279,12 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
       assert((PropVal == 0 || PropVal == 128 || PropVal == 256) &&
              "Unsupported GRF Size");
       // Map sycl-grf-size values to RegisterAllocMode values used in SPIR-V.
+      static constexpr int SMALL_GRF_REGALLOCMODE_VAL = 1;
+      static constexpr int LARGE_GRF_REGALLOCMODE_VAL = 2;
       if (PropVal == 128)
-        PropVal = 1;
+        PropVal = SMALL_GRF_REGALLOCMODE_VAL;
       else if (PropVal == 256)
-        PropVal = 2;
+        PropVal = LARGE_GRF_REGALLOCMODE_VAL;
     }
     Metadata *AttrMDArgs[] = {ConstantAsMetadata::get(
         Constant::getIntegerValue(Type::getInt32Ty(Ctx), APInt(32, PropVal)))};

llvm/test/tools/sycl-post-link/grf-size-conflict.ll

@@ -2,7 +2,7 @@
 
 ; RUN: not sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s 2>&1 | FileCheck %s
 
-; CHECK: Unsupported use of both sycl-register-alloc-mode and sycl-grf-size
+; CHECK: Unsupported use of both register_alloc_mode and grf_size
 
 source_filename = "llvm-link"
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"

llvm/test/tools/sycl-post-link/sycl-grf-size.ll

@@ -2,7 +2,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_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
+; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR --implicit-check-not='__ESIMD_kernel()'
 ; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
 ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-SYCL-LargeGRF-IR
 ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-SYCL-LargeGRF-PROP
@@ -55,7 +55,6 @@ entry:
 }
 
 define weak_odr dso_local spir_kernel void @__ESIMD_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
-
 entry:
   ret void
 }

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -485,8 +485,8 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
     });
     if (HasGRFSize) {
       if (HasRegAllocMode)
-        error("Unsupported use of both sycl-register-alloc-mode and "
-              "sycl-grf-size");
+        error("Unsupported use of both register_alloc_mode and "
+              "grf_size");
       PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({GRFSizeAttr, GRFSizeVal});
     }
   }

sycl/include/sycl/detail/kernel_properties.hpp

@@ -29,9 +29,10 @@ struct register_alloc_mode_key {
 };
 
 template <register_alloc_mode_enum Mode>
-inline constexpr register_alloc_mode_key::value_t<Mode>
-    register_alloc_mode __SYCL_DEPRECATED("register_alloc_mode is deprecated, "
-                                          "use grf_size or grf_size_automatic");
+inline constexpr register_alloc_mode_key::value_t<Mode> register_alloc_mode
+    __SYCL_DEPRECATED("register_alloc_mode is deprecated, "
+                      "use sycl::ext::intel::experimental::grf_size or "
+                      "sycl::ext::intel::experimental::grf_size_automatic");
 } // namespace detail
 
 namespace ext::oneapi::experimental {

sycl/test-e2e/DeviceCodeSplit/grf.cpp

@@ -32,10 +32,11 @@
 // RUN: env SYCL_PROGRAM_COMPILE_OPTIONS="-g" SYCL_PI_TRACE=-1 %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK,CHECK-AUTO-WITH-VAR
 #include "../helpers.hpp"
 #include <iostream>
-#include <sycl/detail/kernel_properties.hpp>
 #include <sycl/sycl.hpp>
 #ifdef USE_NEW_API
 #include <sycl/ext/intel/experimental/grf_size_properties.hpp>
+#else
+#include <sycl/detail/kernel_properties.hpp>
 #endif
 
 using namespace sycl;

sycl/test-e2e/ESIMD/large-grf.cpp

@@ -26,11 +26,12 @@
 #include "esimd_test_utils.hpp"
 
 #include <iostream>
-#include <sycl/detail/kernel_properties.hpp>
 #include <sycl/ext/intel/esimd.hpp>
 #include <sycl/sycl.hpp>
 #ifdef USE_NEW_API
 #include <sycl/ext/intel/experimental/grf_size_properties.hpp>
+#else
+#include <sycl/detail/kernel_properties.hpp>
 #endif
 
 using namespace sycl;