Skip to content

Commit d87afb8

Browse files
sarnexclaude
sarnex
and
claude
authored
[SYCL] Implement sycl_ext_intel_maximum_registers (#22923)
This PR implements [sycl_ext_intel_maximum_registers](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_intel_maximum_registers.asciidoc) which allows per-kernel GRF size specification, communicated to the graphics driver using the [SPV_INTEL_maximum_registers](https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_maximum_registers.asciidoc) SPIR-V extension. The SYCL extension defines two new properties, `maximum_registers<N>` and `maximum_registers_automatic`. There is a similar extension [sycl_ext_intel_grf_size](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_intel_grf_size.asciidoc), however the mechanism by which that extension communicates the user specification to the graphics driver is not supported in some use cases, does not use a first-class SPIR-V extension. The new extension covers all use cases, but it requires a graphics driver supporting [SPV_INTEL_maximum_registers](https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_maximum_registers.asciidoc) and using it will error on drivers that do not support the new extension, so we introduced a new extension to not break existing users of [sycl_ext_intel_grf_size](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_intel_grf_size.asciidoc) and allow users to opt-in to the requirement to have a driver supporting this extension. The property ends up after CFE codegen as the IR attribute `sycl-maximum-registers` which `CompileTimePropertiesPass` processes by adding the `MaximumRegisters` metadata to the kernel, which the SPIR-V translator is looking for. SPIR-V translator support was implemented in KhronosGroup/SPIRV-LLVM-Translator@ca3769b. There is no public driver implementing [SPV_INTEL_maximum_registers](https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_maximum_registers.asciidoc) yet, so I wrote a compile-time test to verify the property mutual exclusivity/supported register sizes, and a currently-build-only E2E test to make sure all cases at least compile. Once a driver is available and we are using it in CI, the test will XPASS so we can enable it. I manually tested this on Linux BMG using a prototype graphics driver and everything worked as expected. I will test on Windows when a final implementation is merged, the prototype didn't support Windows. --------- Signed-off-by: Nick Sarnie <nick.sarnie@intel.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
1 parent 2d1ac7b commit d87afb8

15 files changed

Lines changed: 430 additions & 13 deletions

File tree

‎clang/lib/Driver/ToolChains/Clang.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11288,7 +11288,8 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
1128811288
",+SPV_KHR_cooperative_matrix"
1128911289
",+SPV_EXT_shader_atomic_float16_add"
1129011290
",+SPV_INTEL_fp_max_error"
11291-
",+SPV_INTEL_memory_access_aliasing";
11291+
",+SPV_INTEL_memory_access_aliasing"
11292+
",+SPV_INTEL_maximum_registers";
1129211293

1129311294
TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg));
1129411295
}

‎clang/test/Driver/sycl-spirv-ext-old-model.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,3 +50,4 @@
5050
// CHECK-DEFAULT-SAME:,+SPV_EXT_shader_atomic_float16_add
5151
// CHECK-DEFAULT-SAME:,+SPV_INTEL_fp_max_error
5252
// CHECK-DEFAULT-SAME:,+SPV_INTEL_memory_access_aliasing
53+
// CHECK-DEFAULT-SAME:,+SPV_INTEL_maximum_registers

‎clang/test/Driver/sycl-spirv-ext.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -68,3 +68,4 @@
6868
// CHECK-DEFAULT-SAME:,+SPV_EXT_shader_atomic_float16_add
6969
// CHECK-DEFAULT-SAME:,+SPV_INTEL_fp_max_error
7070
// CHECK-DEFAULT-SAME:,+SPV_INTEL_memory_access_aliasing
71+
// CHECK-DEFAULT-SAME:,+SPV_INTEL_maximum_registers

‎clang/test/Driver/sycl-spirv-metadata-old-model.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
// RUN: FileCheck -check-prefix CHECK-WITHOUT %s
1010

1111
// CHECK-WITH: llvm-spirv{{.*}} "--spirv-preserve-auxdata"
12-
// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
12+
// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_maximum_registers"
1313

1414
// CHECK-WITHOUT: "{{.*}}llvm-spirv"
1515
// CHECK-WITHOUT-NOT: --spirv-preserve-auxdata

‎clang/test/Driver/sycl-spirv-obj-old-model.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111
// SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]"
1212
// SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]"
1313
// SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata"
14-
// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
14+
// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_maximum_registers"
1515
// SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]"
1616
// SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"
1717
// SPIRV_DEVICE_OBJ-SAME: "-fsycl-is-host"

‎clang/test/Driver/sycl-spirv-obj.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111
// SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]"
1212
// SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]"
1313
// SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata"
14-
// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
14+
// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_maximum_registers"
1515
// SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]"
1616
// SPIRV_DEVICE_OBJ: llvm-offload-binary{{.*}} "--image=file=[[DEVICE_SPV]]{{.*}}"
1717
// SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"

‎clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -951,7 +951,8 @@ getTripleBasedSPIRVTransOpts(const ArgList &Args,
951951
",+SPV_KHR_cooperative_matrix"
952952
",+SPV_EXT_shader_atomic_float16_add"
953953
",+SPV_INTEL_fp_max_error"
954-
",+SPV_INTEL_memory_access_aliasing";
954+
",+SPV_INTEL_memory_access_aliasing"
955+
",+SPV_INTEL_maximum_registers";
955956
TranslatorArgs.push_back(Args.MakeArgString(ExtArg));
956957
}
957958

‎llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,7 @@ namespace {
2929
constexpr StringRef SyclHostAccessAttr = "sycl-host-access";
3030
constexpr StringRef SyclPipelinedAttr = "sycl-pipelined";
3131
constexpr StringRef SyclGrfSizeAttr = "sycl-grf-size";
32+
constexpr StringRef SyclMaximumRegistersAttr = "sycl-maximum-registers";
3233

3334
constexpr StringRef SpirvDecorMdKind = "spirv.Decorations";
3435
constexpr StringRef SpirvDecorCacheControlMdKind =
@@ -501,6 +502,21 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
501502
MDNode::get(Ctx, AttrMDArgs));
502503
}
503504

505+
if (AttrKindStr == SyclMaximumRegistersAttr) {
506+
uint32_t PropVal = getAttributeAsInteger<uint32_t>(Attr);
507+
// The property supports only 0, 128, 256 and 512.
508+
if (PropVal != 0 && PropVal != 128 && PropVal != 256 && PropVal != 512)
509+
return std::nullopt;
510+
Metadata *AttrMDArgs[1];
511+
if (PropVal == 0)
512+
AttrMDArgs[0] = MDString::get(Ctx, "AutoINTEL");
513+
else
514+
AttrMDArgs[0] = ConstantAsMetadata::get(
515+
Constant::getIntegerValue(Type::getInt32Ty(Ctx), APInt(32, PropVal)));
516+
return std::pair<std::string, MDNode *>("MaximumRegisters",
517+
MDNode::get(Ctx, AttrMDArgs));
518+
}
519+
504520
return std::nullopt;
505521
}
506522

sycl/doc/extensions/proposed/sycl_ext_intel_maximum_registers.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_intel_maximum_registers.asciidoc

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -50,9 +50,9 @@ This extension also depends on the following other SYCL extensions:
5050

5151
== Status
5252

53-
This is a proposed extension specification, intended to gather community feedback.
54-
Interfaces defined in this specification may not be implemented yet or may be in a preliminary state.
55-
The specification itself may also change in incompatible ways before it is finalized.
53+
This is an experimental extension specification, intended to provide early access to features and gather community feedback.
54+
Interfaces defined in this specification are implemented in DPC++,
55+
but they are not finalized and may change incompatibly in future versions of DPC++ without prior notice.
5656
**Shipping software products should not rely on APIs defined in this specification.**
5757

5858
== Overview

‎sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp‎

Lines changed: 42 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@
88

99
#pragma once
1010

11+
#include <sycl/ext/intel/experimental/maximum_registers_properties.hpp>
1112
#include <sycl/ext/oneapi/free_function_kernel_properties.hpp>
1213
#include <sycl/ext/oneapi/properties.hpp>
1314
#include <sycl/ext/oneapi/properties/property.hpp>
@@ -53,17 +54,54 @@ struct PropertyMetaInfo<
5354
static constexpr unsigned int value = 0;
5455
};
5556

57+
// grf_size, grf_size_automatic, maximum_registers, and
58+
// maximum_registers_automatic are all mutually exclusive.
5659
template <typename Properties>
5760
struct ConflictingProperties<sycl::ext::intel::experimental::grf_size_key,
5861
Properties>
59-
: std::bool_constant<Properties::template has_property<
60-
sycl::ext::intel::experimental::grf_size_automatic_key>()> {};
62+
: std::bool_constant<
63+
Properties::template has_property<
64+
sycl::ext::intel::experimental::grf_size_automatic_key>() ||
65+
Properties::template has_property<
66+
sycl::ext::intel::experimental::maximum_registers_key>() ||
67+
Properties::template has_property<
68+
sycl::ext::intel::experimental::
69+
maximum_registers_automatic_key>()> {};
6170

6271
template <typename Properties>
6372
struct ConflictingProperties<
6473
sycl::ext::intel::experimental::grf_size_automatic_key, Properties>
65-
: std::bool_constant<Properties::template has_property<
66-
sycl::ext::intel::experimental::grf_size_key>()> {};
74+
: std::bool_constant<
75+
Properties::template has_property<
76+
sycl::ext::intel::experimental::grf_size_key>() ||
77+
Properties::template has_property<
78+
sycl::ext::intel::experimental::maximum_registers_key>() ||
79+
Properties::template has_property<
80+
sycl::ext::intel::experimental::
81+
maximum_registers_automatic_key>()> {};
82+
83+
template <typename Properties>
84+
struct ConflictingProperties<
85+
sycl::ext::intel::experimental::maximum_registers_key, Properties>
86+
: std::bool_constant<
87+
Properties::template has_property<
88+
sycl::ext::intel::experimental::grf_size_key>() ||
89+
Properties::template has_property<
90+
sycl::ext::intel::experimental::grf_size_automatic_key>() ||
91+
Properties::template has_property<
92+
sycl::ext::intel::experimental::
93+
maximum_registers_automatic_key>()> {};
94+
95+
template <typename Properties>
96+
struct ConflictingProperties<
97+
sycl::ext::intel::experimental::maximum_registers_automatic_key, Properties>
98+
: std::bool_constant<
99+
Properties::template has_property<
100+
sycl::ext::intel::experimental::grf_size_key>() ||
101+
Properties::template has_property<
102+
sycl::ext::intel::experimental::grf_size_automatic_key>() ||
103+
Properties::template has_property<
104+
sycl::ext::intel::experimental::maximum_registers_key>()> {};
67105

68106
} // namespace ext::oneapi::experimental::detail
69107
} // namespace _V1

0 commit comments

Comments
 (0)