[SYCL] Move GRF property conflict check to compile time (intel#10116)

Using the new property conflict checking infrastructure, we can do the
`sycl::detail::register_alloc_mode` vs
`sycl::experimental::intel::grf_size`/`sycl::experimental::intel::grf_size_automatic`
check at compile time. We already do
`sycl::experimental::intel::grf_size` vs
`sycl::experimental::intel::grf_size_automatic` at compile time.

I can add a test if we think it has value, but my initial reaction is we
would mostly be testing the conflicting checking infrastructure, which
we already have tests for.

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

Author: sarnex

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

@@ -454,12 +454,11 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
   if (MD.isESIMD()) {
     PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true});
   }
-  bool HasRegAllocMode = false;
   {
     StringRef RegAllocModeAttr = "sycl-register-alloc-mode";
     uint32_t RegAllocModeVal;
 
-    HasRegAllocMode = llvm::any_of(MD.entries(), [&](const Function *F) {
+    bool HasRegAllocMode = llvm::any_of(MD.entries(), [&](const Function *F) {
       if (!F->hasFnAttribute(RegAllocModeAttr))
         return false;
       const auto &Attr = F->getFnAttribute(RegAllocModeAttr);
@@ -484,9 +483,6 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
       return true;
     });
     if (HasGRFSize) {
-      if (HasRegAllocMode)
-        error("Unsupported use of both register_alloc_mode and "
-              "grf_size");
       PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({GRFSizeAttr, GRFSizeVal});
     }
   }

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

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/detail/kernel_properties.hpp>
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 
@@ -77,14 +78,30 @@ struct PropertyMetaInfo<
 template <typename Properties>
 struct ConflictingProperties<sycl::ext::intel::experimental::grf_size_key,
                              Properties>
-    : ContainsProperty<sycl::ext::intel::experimental::grf_size_automatic_key,
-                       Properties> {};
+    : std::bool_constant<
+          ContainsProperty<
+              sycl::ext::intel::experimental::grf_size_automatic_key,
+              Properties>::value ||
+          ContainsProperty<sycl::detail::register_alloc_mode_key,
+                           Properties>::value> {};
 
 template <typename Properties>
 struct ConflictingProperties<
     sycl::ext::intel::experimental::grf_size_automatic_key, Properties>
-    : ContainsProperty<sycl::ext::intel::experimental::grf_size_key,
-                       Properties> {};
+    : std::bool_constant<
+          ContainsProperty<sycl::ext::intel::experimental::grf_size_key,
+                           Properties>::value ||
+          ContainsProperty<sycl::detail::register_alloc_mode_key,
+                           Properties>::value> {};
+
+template <typename Properties>
+struct ConflictingProperties<sycl::detail::register_alloc_mode_key, Properties>
+    : std::bool_constant<
+          ContainsProperty<sycl::ext::intel::experimental::grf_size_key,
+                           Properties>::value ||
+          ContainsProperty<
+              sycl::ext::intel::experimental::grf_size_automatic_key,
+              Properties>::value> {};
 
 } // namespace detail
 } // namespace ext::oneapi::experimental