Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion clang/lib/Driver/ToolChains/Clang.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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";

@sarnex sarnex Aug 11, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This will not cause an issue for older graphics drivers. The only way to generate a SPIR-V module that actually uses this extension is to use the new SYCL property implemented in this PR.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ah good catch, i incorrectly assumed the driver one was used for both. fixed in b403905, thanks


TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg));
}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/Driver/sycl-spirv-metadata-old-model.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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
2 changes: 1 addition & 1 deletion clang/test/Driver/sycl-spirv-obj-old-model.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion clang/test/Driver/sycl-spirv-obj.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down
3 changes: 2 additions & 1 deletion clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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));
}

Expand Down
16 changes: 16 additions & 0 deletions llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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 =
Expand Down Expand Up @@ -501,6 +502,21 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
MDNode::get(Ctx, AttrMDArgs));
}

if (AttrKindStr == SyclMaximumRegistersAttr) {
uint32_t PropVal = getAttributeAsInteger<uint32_t>(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<std::string, MDNode *>("MaximumRegisters",
MDNode::get(Ctx, AttrMDArgs));
}

return std::nullopt;
}

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@

#pragma once

#include <sycl/ext/intel/experimental/maximum_registers_properties.hpp>
#include <sycl/ext/oneapi/free_function_kernel_properties.hpp>
#include <sycl/ext/oneapi/properties.hpp>
#include <sycl/ext/oneapi/properties/property.hpp>
Expand Down Expand Up @@ -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 <typename Properties>
struct ConflictingProperties<sycl::ext::intel::experimental::grf_size_key,
Properties>
: std::bool_constant<Properties::template has_property<
sycl::ext::intel::experimental::grf_size_automatic_key>()> {};
: 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 <typename Properties>
struct ConflictingProperties<
sycl::ext::intel::experimental::grf_size_automatic_key, Properties>
: std::bool_constant<Properties::template has_property<
sycl::ext::intel::experimental::grf_size_key>()> {};
: 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 <typename Properties>
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 <typename Properties>
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
Expand Down
Original file line number Diff line number Diff line change
@@ -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 <sycl/ext/oneapi/free_function_kernel_properties.hpp>
#include <sycl/ext/oneapi/properties.hpp>
#include <sycl/ext/oneapi/properties/property.hpp>
#include <sycl/ext/oneapi/properties/property_value.hpp>

#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 <unsigned int Size>
using value_t = oneapi::experimental::property_value<
maximum_registers_key, std::integral_constant<unsigned int, Size>>;
};

struct maximum_registers_automatic_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::MaximumRegistersAutomatic> {
using value_t =
oneapi::experimental::property_value<maximum_registers_automatic_key>;
};

template <unsigned int Size>
inline constexpr maximum_registers_key::value_t<Size> maximum_registers;

inline constexpr maximum_registers_automatic_key::value_t
maximum_registers_automatic;

} // namespace ext::intel::experimental
namespace ext::oneapi::experimental::detail {
template <unsigned int Size>
struct PropertyMetaInfo<
sycl::ext::intel::experimental::maximum_registers_key::value_t<Size>> {
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
4 changes: 3 additions & 1 deletion sycl/include/sycl/ext/oneapi/properties/property.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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 <typename PropertyT> struct PropertyToKind {
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/sycl.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -101,6 +101,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")

#include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp>
#include <sycl/ext/intel/experimental/grf_size_properties.hpp>
#include <sycl/ext/intel/experimental/maximum_registers_properties.hpp>
#include <sycl/ext/intel/experimental/usm_properties.hpp>
#include <sycl/ext/oneapi/bfloat16.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
Expand Down
Loading
Loading