Skip to content

[SYCL][ABI-break] Remove deprecated register_alloc_mode kernel property #20919

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
sarnex merged 1 commit into from
Dec 18, 2025
Merged

[SYCL][ABI-break] Remove deprecated register_alloc_mode kernel property #20919

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
31 changes: 13 additions & 18 deletions llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,6 @@ namespace {

constexpr StringRef SyclHostAccessAttr = "sycl-host-access";
constexpr StringRef SyclPipelinedAttr = "sycl-pipelined";
constexpr StringRef SyclRegisterAllocModeAttr = "sycl-register-alloc-mode";
constexpr StringRef SyclGrfSizeAttr = "sycl-grf-size";

constexpr StringRef SpirvDecorMdKind = "spirv.Decorations";
Expand Down Expand Up @@ -482,24 +481,20 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
MDNode::get(Ctx, ClusterMDArgs));
}

if ((AttrKindStr == SyclRegisterAllocModeAttr ||
AttrKindStr == SyclGrfSizeAttr) &&
!llvm::esimd::isESIMD(F)) {
// TODO: Remove SYCL_REGISTER_ALLOC_MODE_ATTR support in next ABI break.
if ((AttrKindStr == SyclGrfSizeAttr) && !llvm::esimd::isESIMD(F)) {
uint32_t PropVal = getAttributeAsInteger<uint32_t>(Attr);
if (AttrKindStr == SyclGrfSizeAttr) {
// The RegisterAllocMode metadata supports only 0, 128, and 256 for
// PropVal.
if (PropVal != 0 && PropVal != 128 && PropVal != 256)
return std::nullopt;
// 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 = SMALL_GRF_REGALLOCMODE_VAL;
else if (PropVal == 256)
PropVal = LARGE_GRF_REGALLOCMODE_VAL;
}
// The RegisterAllocMode metadata supports only 0, 128, and 256 for
// PropVal.
if (PropVal != 0 && PropVal != 128 && PropVal != 256)
return std::nullopt;
// 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 = SMALL_GRF_REGALLOCMODE_VAL;
else if (PropVal == 256)
PropVal = LARGE_GRF_REGALLOCMODE_VAL;

Metadata *AttrMDArgs[] = {ConstantAsMetadata::get(
Constant::getIntegerValue(Type::getInt32Ty(Ctx), APInt(32, PropVal)))};
return std::pair<std::string, MDNode *>("RegisterAllocMode",
Expand Down
17 changes: 0 additions & 17 deletions llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -298,23 +298,6 @@ PropSetRegTy computeModuleProperties(const Module &M,

if (SplitType == module_split::SyclEsimdSplitStatus::ESIMD_ONLY)
PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "isEsimdImage", true);
{
StringRef RegAllocModeAttr = "sycl-register-alloc-mode";
uint32_t RegAllocModeVal;

bool HasRegAllocMode = llvm::any_of(EntryPoints, [&](const Function *F) {
if (!F->hasFnAttribute(RegAllocModeAttr))
return false;
const auto &Attr = F->getFnAttribute(RegAllocModeAttr);
RegAllocModeVal = getAttributeAsInteger<uint32_t>(Attr);
return true;
});
if (HasRegAllocMode) {
PropSet.add(PropSetRegTy::SYCL_MISC_PROP, RegAllocModeAttr,
RegAllocModeVal);
}
}

{
StringRef GRFSizeAttr = "sycl-grf-size";
uint32_t GRFSizeVal;
Expand Down
1 change: 0 additions & 1 deletion llvm/lib/SYCLPostLink/ModuleSplitter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1002,7 +1002,6 @@ std::string computeFuncCategoryForSplittingPerSource(const Function &F) {
// Optional features
// NOTE: Add more categories at the end of the list to avoid changing orders
// of output files in existing tests.
computeFuncCategoryFromAttribute(F, "sycl-register-alloc-mode", Result);
computeFuncCategoryFromAttribute(F, "sycl-grf-size", Result);
computeFuncCategoryFromSYCLUsedAspects(F, Result);
computeFuncCategoryFromIntegersListMetadata(F, "reqd_work_group_size",
Expand Down
35 changes: 0 additions & 35 deletions llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-alloc-mode.ll
View file
Open in desktop

This file was deleted.

81 changes: 0 additions & 81 deletions llvm/test/tools/sycl-post-link/registerallocmode.ll
View file
Open in desktop

This file was deleted.

1 change: 0 additions & 1 deletion sycl/doc/design/PropertySets.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -109,7 +109,6 @@ Miscellaneous properties:
| Key | Value type | Value |
| ------------------------------- | --------------------- | ---------------------------------------------------------------------------------------------------------------- |
| "isEsimdImage" | 32 bit integer. ("1") | 1 if the image is ESIMD and 0 or missing otherwise. |
| "sycl-register-alloc-mode" | 32 bit integer. ("1") | The register allocation mode: 2 for large and 0 or missing for automatic. |
| "sycl-grf-size" | 32 bit integer. ("1") | The GRF size. Automatic if 0 or missing. |
| "optLevel" | 32 bit integer. ("1") | Optimization level, corresponding to the `-O` option used during compilation. |
| "sanUsed" | Byte array. ("2") | Specifying if address sanitization ("asan") or memory sanitization ("msan") is used. Missing if neither is used. |
Expand Down
51 changes: 0 additions & 51 deletions sycl/include/sycl/detail/kernel_properties.hpp
View file
Open in desktop

This file was deleted.

20 changes: 3 additions & 17 deletions sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,6 @@

#pragma once

#include <sycl/detail/kernel_properties.hpp>
#include <sycl/ext/oneapi/properties/properties.hpp>
#include <sycl/ext/oneapi/properties/property.hpp>
#include <sycl/ext/oneapi/properties/property_value.hpp>
Expand Down Expand Up @@ -55,27 +54,14 @@ struct PropertyMetaInfo<
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>() ||
Properties::template has_property<
sycl::detail::register_alloc_mode_key>()> {};
: std::bool_constant<Properties::template has_property<
sycl::ext::intel::experimental::grf_size_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>() ||
Properties::template has_property<
sycl::detail::register_alloc_mode_key>()> {};

template <typename Properties>
struct ConflictingProperties<sycl::detail::register_alloc_mode_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>()> {};
sycl::ext::intel::experimental::grf_size_key>()> {};

} // namespace ext::oneapi::experimental::detail
} // namespace _V1
Expand Down
85 changes: 42 additions & 43 deletions sycl/include/sycl/ext/oneapi/properties/property.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -151,50 +151,49 @@ enum PropKind : uint32_t {
Alignment = 6,
CacheConfig = 7,
UseRootSync = 8,
RegisterAllocMode = 9,
GRFSize = 10,
GRFSizeAutomatic = 11,
ESIMDL1CacheHint = 12,
ESIMDL2CacheHint = 13,
ESIMDL3CacheHint = 14,
UsmKind = 15,
CacheControlReadHint = 16,
CacheControlReadAssertion = 17,
CacheControlWrite = 18,
BuildOptions = 19,
BuildLog = 20,
FloatingPointControls = 21,
DataPlacement = 22,
ContiguousMemory = 23,
FullGroup = 24,
Naive = 25,
WorkGroupProgress = 26,
SubGroupProgress = 27,
WorkItemProgress = 28,
NDRangeKernel = 29,
SingleTaskKernel = 30,
IndirectlyCallable = 31,
CallsIndirectly = 32,
InputDataPlacement = 33,
OutputDataPlacement = 34,
IncludeFiles = 35,
RegisteredNames = 36,
ClusterLaunch = 37,
MaxWorkGroupSize = 38,
MaxLinearWorkGroupSize = 39,
Prefetch = 40,
Deterministic = 41,
InitializeToIdentity = 42,
WorkGroupScratchSize = 43,
Unaliased = 44,
EventMode = 45,
NativeLocalBlockIO = 46,
InitialThreshold = 47,
MaximumSize = 48,
ZeroInit = 49,
FastLink = 50,
GRFSize = 9,
Copy link
Contributor

@cperkinsintel cperkinsintel Dec 17, 2025

Choose a reason for hiding this comment

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

can we remove enum values? Won't that result in ABI breaks?

I guess "tis the season" ??

Copy link
Contributor Author

@againull againull Dec 17, 2025

Choose a reason for hiding this comment

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

I removed the RegisterAllocMode. But GRFSize ones are in use.

GRFSizeAutomatic = 10,
ESIMDL1CacheHint = 11,
ESIMDL2CacheHint = 12,
ESIMDL3CacheHint = 13,
UsmKind = 14,
CacheControlReadHint = 15,
CacheControlReadAssertion = 16,
CacheControlWrite = 17,
BuildOptions = 18,
BuildLog = 19,
FloatingPointControls = 20,
DataPlacement = 21,
ContiguousMemory = 22,
FullGroup = 23,
Naive = 24,
WorkGroupProgress = 25,
SubGroupProgress = 26,
WorkItemProgress = 27,
NDRangeKernel = 28,
SingleTaskKernel = 29,
IndirectlyCallable = 30,
CallsIndirectly = 31,
InputDataPlacement = 32,
OutputDataPlacement = 33,
IncludeFiles = 34,
RegisteredNames = 35,
ClusterLaunch = 36,
MaxWorkGroupSize = 37,
MaxLinearWorkGroupSize = 38,
Prefetch = 39,
Deterministic = 40,
InitializeToIdentity = 41,
WorkGroupScratchSize = 42,
Unaliased = 43,
EventMode = 44,
NativeLocalBlockIO = 45,
InitialThreshold = 46,
MaximumSize = 47,
ZeroInit = 48,
FastLink = 49,
// PropKindSize must always be the last value.
PropKindSize = 51,
PropKindSize = 50,
};

template <typename PropertyT> struct PropertyToKind {
Expand Down
Loading