diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 0aa88c190ad0f..cab1f970613cc 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11288,7 +11288,8 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C, ",+SPV_KHR_cooperative_matrix" ",+SPV_EXT_shader_atomic_float16_add" ",+SPV_INTEL_fp_max_error" - ",+SPV_INTEL_memory_access_aliasing"; + ",+SPV_INTEL_memory_access_aliasing" + ",+SPV_INTEL_maximum_registers"; TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg)); } diff --git a/clang/test/Driver/sycl-spirv-metadata-old-model.cpp b/clang/test/Driver/sycl-spirv-metadata-old-model.cpp index 2e19dd9ed1dc4..81549fab08d2e 100644 --- a/clang/test/Driver/sycl-spirv-metadata-old-model.cpp +++ b/clang/test/Driver/sycl-spirv-metadata-old-model.cpp @@ -9,7 +9,7 @@ // RUN: FileCheck -check-prefix CHECK-WITHOUT %s // CHECK-WITH: llvm-spirv{{.*}} "--spirv-preserve-auxdata" -// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing" +// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_maximum_registers" // CHECK-WITHOUT: "{{.*}}llvm-spirv" // CHECK-WITHOUT-NOT: --spirv-preserve-auxdata diff --git a/clang/test/Driver/sycl-spirv-obj-old-model.cpp b/clang/test/Driver/sycl-spirv-obj-old-model.cpp index 7a5f6215e31ae..da8bdacb97d58 100644 --- a/clang/test/Driver/sycl-spirv-obj-old-model.cpp +++ b/clang/test/Driver/sycl-spirv-obj-old-model.cpp @@ -11,7 +11,7 @@ // SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]" // SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]" // SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata" -// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing" +// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_maximum_registers" // SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]" // SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu" // SPIRV_DEVICE_OBJ-SAME: "-fsycl-is-host" diff --git a/clang/test/Driver/sycl-spirv-obj.cpp b/clang/test/Driver/sycl-spirv-obj.cpp index 1d3aaf2803c88..dd5deb8a2da69 100644 --- a/clang/test/Driver/sycl-spirv-obj.cpp +++ b/clang/test/Driver/sycl-spirv-obj.cpp @@ -11,7 +11,7 @@ // SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]" // SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]" // SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata" -// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing" +// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_maximum_registers" // SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]" // SPIRV_DEVICE_OBJ: llvm-offload-binary{{.*}} "--image=file=[[DEVICE_SPV]]{{.*}}" // SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu" diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 59aa81ed94a74..8a846aa341b0b 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -951,7 +951,8 @@ getTripleBasedSPIRVTransOpts(const ArgList &Args, ",+SPV_KHR_cooperative_matrix" ",+SPV_EXT_shader_atomic_float16_add" ",+SPV_INTEL_fp_max_error" - ",+SPV_INTEL_memory_access_aliasing"; + ",+SPV_INTEL_memory_access_aliasing" + ",+SPV_INTEL_maximum_registers"; TranslatorArgs.push_back(Args.MakeArgString(ExtArg)); } diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 48c64e2d41976..313bbdc46f7df 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -29,6 +29,7 @@ namespace { constexpr StringRef SyclHostAccessAttr = "sycl-host-access"; constexpr StringRef SyclPipelinedAttr = "sycl-pipelined"; constexpr StringRef SyclGrfSizeAttr = "sycl-grf-size"; +constexpr StringRef SyclMaximumRegistersAttr = "sycl-maximum-registers"; constexpr StringRef SpirvDecorMdKind = "spirv.Decorations"; constexpr StringRef SpirvDecorCacheControlMdKind = @@ -501,6 +502,21 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) { MDNode::get(Ctx, AttrMDArgs)); } + if (AttrKindStr == SyclMaximumRegistersAttr) { + uint32_t PropVal = getAttributeAsInteger(Attr); + // The property supports only 0, 128, 256 and 512. + if (PropVal != 0 && PropVal != 128 && PropVal != 256 && PropVal != 512) + return std::nullopt; + Metadata *AttrMDArgs[1]; + if (PropVal == 0) + AttrMDArgs[0] = MDString::get(Ctx, "AutoINTEL"); + else + AttrMDArgs[0] = ConstantAsMetadata::get( + Constant::getIntegerValue(Type::getInt32Ty(Ctx), APInt(32, PropVal))); + return std::pair("MaximumRegisters", + MDNode::get(Ctx, AttrMDArgs)); + } + return std::nullopt; } diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_maximum_registers.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_intel_maximum_registers.asciidoc similarity index 96% rename from sycl/doc/extensions/proposed/sycl_ext_intel_maximum_registers.asciidoc rename to sycl/doc/extensions/experimental/sycl_ext_intel_maximum_registers.asciidoc index aee38e9601aa2..b43c1f587b44b 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_maximum_registers.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_intel_maximum_registers.asciidoc @@ -50,9 +50,9 @@ This extension also depends on the following other SYCL extensions: == Status -This is a proposed extension specification, intended to gather community feedback. -Interfaces defined in this specification may not be implemented yet or may be in a preliminary state. -The specification itself may also change in incompatible ways before it is finalized. +This is an experimental extension specification, intended to provide early access to features and gather community feedback. +Interfaces defined in this specification are implemented in DPC++, +but they are not finalized and may change incompatibly in future versions of DPC++ without prior notice. **Shipping software products should not rely on APIs defined in this specification.** == Overview diff --git a/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp b/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp index 395555596fcd6..9a57ee021e96b 100644 --- a/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp @@ -8,6 +8,7 @@ #pragma once +#include #include #include #include @@ -53,17 +54,54 @@ struct PropertyMetaInfo< static constexpr unsigned int value = 0; }; +// grf_size, grf_size_automatic, maximum_registers, and +// maximum_registers_automatic are all mutually exclusive. template struct ConflictingProperties - : std::bool_constant()> {}; + : std::bool_constant< + Properties::template has_property< + sycl::ext::intel::experimental::grf_size_automatic_key>() || + Properties::template has_property< + sycl::ext::intel::experimental::maximum_registers_key>() || + Properties::template has_property< + sycl::ext::intel::experimental:: + maximum_registers_automatic_key>()> {}; template struct ConflictingProperties< sycl::ext::intel::experimental::grf_size_automatic_key, Properties> - : std::bool_constant()> {}; + : std::bool_constant< + Properties::template has_property< + sycl::ext::intel::experimental::grf_size_key>() || + Properties::template has_property< + sycl::ext::intel::experimental::maximum_registers_key>() || + Properties::template has_property< + sycl::ext::intel::experimental:: + maximum_registers_automatic_key>()> {}; + +template +struct ConflictingProperties< + sycl::ext::intel::experimental::maximum_registers_key, Properties> + : std::bool_constant< + Properties::template has_property< + sycl::ext::intel::experimental::grf_size_key>() || + Properties::template has_property< + sycl::ext::intel::experimental::grf_size_automatic_key>() || + Properties::template has_property< + sycl::ext::intel::experimental:: + maximum_registers_automatic_key>()> {}; + +template +struct ConflictingProperties< + sycl::ext::intel::experimental::maximum_registers_automatic_key, Properties> + : std::bool_constant< + Properties::template has_property< + sycl::ext::intel::experimental::grf_size_key>() || + Properties::template has_property< + sycl::ext::intel::experimental::grf_size_automatic_key>() || + Properties::template has_property< + sycl::ext::intel::experimental::maximum_registers_key>()> {}; } // namespace ext::oneapi::experimental::detail } // namespace _V1 diff --git a/sycl/include/sycl/ext/intel/experimental/maximum_registers_properties.hpp b/sycl/include/sycl/ext/intel/experimental/maximum_registers_properties.hpp new file mode 100644 index 0000000000000..af63f9682719d --- /dev/null +++ b/sycl/include/sycl/ext/intel/experimental/maximum_registers_properties.hpp @@ -0,0 +1,61 @@ +//==---------------- maximum_registers_properties.hpp --------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===--------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include +#include + +#define SYCL_EXT_INTEL_MAXIMUM_REGISTERS 1 + +namespace sycl { +inline namespace _V1 { +namespace ext::intel::experimental { +struct maximum_registers_key + : oneapi::experimental::detail::compile_time_property_key< + oneapi::experimental::detail::PropKind::MaximumRegisters> { + template + using value_t = oneapi::experimental::property_value< + maximum_registers_key, std::integral_constant>; +}; + +struct maximum_registers_automatic_key + : oneapi::experimental::detail::compile_time_property_key< + oneapi::experimental::detail::PropKind::MaximumRegistersAutomatic> { + using value_t = + oneapi::experimental::property_value; +}; + +template +inline constexpr maximum_registers_key::value_t maximum_registers; + +inline constexpr maximum_registers_automatic_key::value_t + maximum_registers_automatic; + +} // namespace ext::intel::experimental +namespace ext::oneapi::experimental::detail { +template +struct PropertyMetaInfo< + sycl::ext::intel::experimental::maximum_registers_key::value_t> { + static_assert(Size == 128 || Size == 256 || Size == 512, + "Unsupported maximum registers"); + static constexpr const char *name = "sycl-maximum-registers"; + static constexpr unsigned int value = Size; +}; +template <> +struct PropertyMetaInfo< + sycl::ext::intel::experimental::maximum_registers_automatic_key::value_t> { + static constexpr const char *name = "sycl-maximum-registers"; + static constexpr unsigned int value = 0; +}; + +} // namespace ext::oneapi::experimental::detail +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index e410173a0f690..79eeff77841ba 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -195,8 +195,10 @@ enum PropKind : uint32_t { EnableIPC = 50, RegisterHostMemoryReadOnly = 51, EnableProfiling = 52, + MaximumRegisters = 53, + MaximumRegistersAutomatic = 54, // PropKindSize must always be the last value. - PropKindSize = 53, + PropKindSize = 55, }; template struct PropertyToKind { diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index d133824184470..88b9065f2020b 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -101,6 +101,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.") #include #include +#include #include #include #include diff --git a/sycl/test-e2e/Properties/maximum_registers.cpp b/sycl/test-e2e/Properties/maximum_registers.cpp new file mode 100644 index 0000000000000..1959184295c32 --- /dev/null +++ b/sycl/test-e2e/Properties/maximum_registers.cpp @@ -0,0 +1,238 @@ +//==----------- maximum_registers.cpp - DPC++ SYCL on-device test --------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// This test verifies the effect of the "maximum_registers" and +// "maximum_registers_automatic" kernel properties from +// sycl_ext_intel_maximum_registers in device code and checks the computed +// results on the host. It exercises the property across these dimensions: +// * SYCL and ESIMD kernels. +// * Free function kernels and lambda kernels. +// * maximum_registers and maximum_registers_automatic. + +// REQUIRES: arch-intel_gpu_bmg_g21 + +// XFAIL: run-mode +// XFAIL-TRACKER: GSD-4149 + +// UNSUPPORTED: spirv-backend +// UNSUPPORTED-INTENDED: The required SPIR-V extensions are not supported. + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// TODO: Add AOT test when removing XFAIL + +#include "../helpers.hpp" +#include +#include +#include +#include +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; +using namespace sycl::ext::intel::experimental; +namespace syclexp = sycl::ext::oneapi::experimental; +namespace intelexp = sycl::ext::intel::experimental; + +constexpr unsigned Size = 32; +constexpr unsigned VL = 16; + +bool checkResult(const std::vector &A, int Inc, const char *Msg) { + int err_cnt = 0; + unsigned Sz = A.size(); + + for (unsigned i = 0; i < Sz; ++i) { + if (A[i] != i + Inc) + if (++err_cnt < 10) + std::cerr << "failed at A[" << i << "]: " << A[i] << " != " << i + Inc + << "\n"; + } + + if (err_cnt > 0) { + std::cout << Msg << " failed. pass rate: " + << ((float)(Sz - err_cnt) / (float)Sz) * 100.0f << "% (" + << (Sz - err_cnt) << "/" << Sz << ")\n"; + return false; + } + std::cout << Msg << " passed\n"; + return true; +} + +// Free function kernels. The property is attached at the definition via +// SYCL_EXT_ONEAPI_FUNCTION_PROPERTY. +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((intelexp::maximum_registers<256>)) +void free_function_kernel_specified(float *Ptr) { + size_t i = ext::oneapi::this_work_item::get_nd_item<1>().get_global_id(0); + Ptr[i] += 1; +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((intelexp::maximum_registers_automatic)) +void free_function_kernel_automatic(float *Ptr) { + size_t i = ext::oneapi::this_work_item::get_nd_item<1>().get_global_id(0); + Ptr[i] += 1; +} + +// ESIMD free function kernels. The property is attached at the definition and +// the function is marked as an ESIMD kernel. +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((intelexp::maximum_registers<256>)) +void esimd_free_function_kernel_specified(float *Ptr) SYCL_ESIMD_KERNEL { + size_t i = ext::oneapi::this_work_item::get_nd_item<1>().get_global_id(0); + float *Base = Ptr + i * VL; + simd va; + va.copy_from(Base); + simd vc = va + 1; + vc.copy_to(Base); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((intelexp::maximum_registers_automatic)) +void esimd_free_function_kernel_automatic(float *Ptr) SYCL_ESIMD_KERNEL { + size_t i = ext::oneapi::this_work_item::get_nd_item<1>().get_global_id(0); + float *Base = Ptr + i * VL; + simd va; + va.copy_from(Base); + simd vc = va + 1; + vc.copy_to(Base); +} + +// SYCL lambda kernel. The property is attached via a launch_config. +template +bool runLambdaSYCL(queue &q, PropsT Props, const char *Msg) { + std::vector A(Size); + for (unsigned i = 0; i < Size; ++i) + A[i] = i; + float *Ptr = malloc_shared(Size, q); + try { + for (unsigned i = 0; i < Size; ++i) + Ptr[i] = A[i]; + + syclexp::parallel_for(q, syclexp::launch_config{range<1>{Size}, Props}, + [=](id<1> i) { Ptr[i] += 1; }); + q.wait(); + + for (unsigned i = 0; i < Size; ++i) + A[i] = Ptr[i]; + } catch (sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + free(Ptr, q); + return false; + } + free(Ptr, q); + return checkResult(A, 1, Msg); +} + +// ESIMD lambda kernel. The property is attached via a launch_config. +template +bool runLambdaESIMD(queue &q, PropsT Props, const char *Msg) { + std::vector A(Size); + for (unsigned i = 0; i < Size; ++i) + A[i] = i; + float *Ptr = malloc_shared(Size, q); + try { + for (unsigned i = 0; i < Size; ++i) + Ptr[i] = A[i]; + + syclexp::parallel_for(q, syclexp::launch_config{range<1>{Size / VL}, Props}, + [=](id<1> i) SYCL_ESIMD_KERNEL { + float *Base = Ptr + i * VL; + simd va; + va.copy_from(Base); + simd vc = va + 1; + vc.copy_to(Base); + }); + q.wait(); + + for (unsigned i = 0; i < Size; ++i) + A[i] = Ptr[i]; + } catch (sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + free(Ptr, q); + return false; + } + free(Ptr, q); + return checkResult(A, 1, Msg); +} + +// Free function kernel. The property is attached at the definition. GlobalSize +// is the number of work-items to launch (Size for SYCL, Size / VL for ESIMD, +// where each work-item processes a VL-wide block). +template +bool runFreeFunction(queue &q, unsigned GlobalSize, const char *Msg) { + std::vector A(Size); + for (unsigned i = 0; i < Size; ++i) + A[i] = i; + float *Ptr = malloc_shared(Size, q); + try { + for (unsigned i = 0; i < Size; ++i) + Ptr[i] = A[i]; + + syclexp::nd_launch(q, + nd_range<1>{range<1>{GlobalSize}, range<1>{GlobalSize}}, + syclexp::kernel_function, Ptr); + q.wait(); + + for (unsigned i = 0; i < Size; ++i) + A[i] = Ptr[i]; + } catch (sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + free(Ptr, q); + return false; + } + free(Ptr, q); + return checkResult(A, 1, Msg); +} + +int main(void) { + queue q(sycl::gpu_selector_v, exceptionHandlerHelper); + + auto dev = q.get_device(); + std::cout << "Running on " << dev.get_info() + << "\n"; + + syclexp::properties specified_props{maximum_registers<256>}; + syclexp::properties automatic_props{maximum_registers_automatic}; + + bool Pass = true; + + // SYCL lambda kernels. + Pass &= + runLambdaSYCL(q, specified_props, "SYCL lambda maximum_registers<256>"); + Pass &= runLambdaSYCL(q, automatic_props, + "SYCL lambda maximum_registers_automatic"); + + // ESIMD lambda kernels. + Pass &= + runLambdaESIMD(q, specified_props, "ESIMD lambda maximum_registers<256>"); + Pass &= runLambdaESIMD(q, automatic_props, + "ESIMD lambda maximum_registers_automatic"); + + // SYCL free function kernels (property attached at the definition). + Pass &= runFreeFunction( + q, Size, "SYCL free function maximum_registers<256>"); + Pass &= runFreeFunction( + q, Size, "SYCL free function maximum_registers_automatic"); + + // ESIMD free function kernels (property attached at the definition). + Pass &= runFreeFunction( + q, Size / VL, "ESIMD free function maximum_registers<256>"); + Pass &= runFreeFunction( + q, Size / VL, "ESIMD free function maximum_registers_automatic"); + + if (!Pass) { + std::cout << "Test failed\n"; + return 1; + } + + std::cout << "Test passed\n"; + return 0; +} diff --git a/sycl/test/extensions/properties/maximum_registers.cpp b/sycl/test/extensions/properties/maximum_registers.cpp new file mode 100644 index 0000000000000..f94d2c3eb4c43 --- /dev/null +++ b/sycl/test/extensions/properties/maximum_registers.cpp @@ -0,0 +1,57 @@ +// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -fsyntax-only -ferror-limit=0 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s + +#include + +struct Kernel { + void operator()(sycl::id<1> i) const {} + + auto get(sycl::ext::oneapi::experimental::properties_tag) const { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::maximum_registers<768>}; + } +}; + +int main() { + // expected-error-re@sycl/ext/oneapi/properties.hpp:* {{static assertion failed due to requirement {{.+}}: Conflicting properties in property list.}} + // expected-note-re@+1 {{in instantiation of function template specialization {{.+}}}} + sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::grf_size<128>, + sycl::ext::intel::experimental::maximum_registers<128>}; + + // expected-error-re@sycl/ext/oneapi/properties.hpp:* {{static assertion failed due to requirement {{.+}}: Conflicting properties in property list.}} + // expected-note-re@+1 {{in instantiation of function template specialization {{.+}}}} + sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::grf_size<128>, + sycl::ext::intel::experimental::grf_size_automatic}; + + // expected-error-re@sycl/ext/oneapi/properties.hpp:* {{static assertion failed due to requirement {{.+}}: Conflicting properties in property list.}} + // expected-note-re@+1 {{in instantiation of function template specialization {{.+}}}} + sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::grf_size<128>, + sycl::ext::intel::experimental::maximum_registers_automatic}; + + // expected-error-re@sycl/ext/oneapi/properties.hpp:* {{static assertion failed due to requirement {{.+}}: Conflicting properties in property list.}} + // expected-note-re@+1 {{in instantiation of function template specialization {{.+}}}} + sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::grf_size_automatic, + sycl::ext::intel::experimental::maximum_registers<128>}; + + // expected-error-re@sycl/ext/oneapi/properties.hpp:* {{static assertion failed due to requirement {{.+}}: Conflicting properties in property list.}} + // expected-note-re@+1 {{in instantiation of function template specialization {{.+}}}} + sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::grf_size_automatic, + sycl::ext::intel::experimental::maximum_registers_automatic}; + + // expected-error-re@sycl/ext/oneapi/properties.hpp:* {{static assertion failed due to requirement {{.+}}: Conflicting properties in property list.}} + // expected-note-re@+1 {{in instantiation of function template specialization {{.+}}}} + sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::maximum_registers<128>, + sycl::ext::intel::experimental::maximum_registers_automatic}; + + sycl::queue q; + // expected-error-re@sycl/ext/intel/experimental/maximum_registers_properties.hpp:* {{static assertion failed due to requirement {{.+}}: Unsupported maximum registers}} + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<1>(1), Kernel{}); + }); + return 0; +}