Skip to content

[SYCL] Remove sycl_ext_oneapi_sub_group in preview-breaking mode - #22883

Open
softwaresat wants to merge 3 commits into
intel:syclfrom
softwaresat:remove-oneapi-sub-group
Open

[SYCL] Remove sycl_ext_oneapi_sub_group in preview-breaking mode#22883
softwaresat wants to merge 3 commits into
intel:syclfrom
softwaresat:remove-oneapi-sub-group

Conversation

@softwaresat

@softwaresat softwaresat commented Aug 6, 2026

Copy link
Copy Markdown
Contributor

Guards the deprecated SYCL_INTEL_sub_group feature-test macro from preview-breaking builds.

Documentation is marked for removal or relocation during the next breaking-changes window.

This closes #22581.

@softwaresat
softwaresat requested review from a team as code owners August 6, 2026 01:48
@softwaresat
softwaresat requested a review from againull August 6, 2026 01:48

@KornevNikita KornevNikita left a comment

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.

This descriptor and all related code should be removed:

sub_group_independent_forward_progress

This header should be removed:
https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/ext/oneapi/sub_group.hpp

I guess the rest of code that should be removed will be easy to find if you disable this header and try to compile the project.

@gmlueck

gmlueck commented Aug 6, 2026

Copy link
Copy Markdown
Contributor

This deprecated extension also defines a couple of C++ attributes:

  • [[intel::sub_group_size(S)]]
  • [[intel::named_sub_group_size(NAME)]]

These are implemented in the C++ front-end (not the headers). I think we should deprecate these also. I don't know if there is a way to remove them from the CFE in -fpreview-breaking-changes. If you want to assign this part of the task to the CFE team, that's fine. But, we should create a tracker in that case.

Note that we have functionality similar to [[intel::named_sub_group_size(NAME)]] proposed in sycl_ext_oneapi_named_sub_group_sizes, but that proposal uses properties (not attributes). Therefore, even if we implement that extension, I don't think we would want that attribute.

@KornevNikita

Copy link
Copy Markdown
Contributor

@softwaresat hi, do you plan to continue working on this?

@softwaresat

Copy link
Copy Markdown
Contributor Author

@KornevNikita Sorry, was on a long vacation and just started classes. I'll have it completed before the end of the week. I'd love to start on other issues afterwards as well.

@softwaresat
softwaresat requested review from a team as code owners August 26, 2026 04:12
|+uint32_t+
|Returns the maximum number of sub-groups in a work-group for any kernel executed on the device. The minimum value is 1.

|+info::device::sub_group_independent_forward_progress+

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.

No need to remove this.

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.

Updated code

struct sub_group;
namespace ext {
namespace oneapi {
struct sub_group;

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.

Please mark this and all other removals with #ifndef __INTEL_PREVIEW_BREAKING_CHANGES. I guess it won't break ABI actually, but I'd like to keep all the oneapi::sub_group related stuff until the next major release.

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.

Fixed all the removals and kept the other sub_group stuff.


void SemaSYCL::checkDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
StringRef NewName) {
// The Intel named sub-group size spelling is deprecated in favor of the

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.

Why do we need this?

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.

I added this to emit a deprecation warning for [[intel::named_sub_group_size(NAME)]] while keeping the attribute source-compatible. Should I remove it?

| Query | CUDA backend result |
| --------------- | ------------------------- |
| `info::device::max_num_sub_groups` | sm 3.0 to 7.0: 64; sm 7.5 32 (see [HW_spec]) |
| `info::device::sub_group_independent_forward_progress` | `true` |

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.

I think we should keep it until the spec is not removed

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.

Added it back

@elizabethandrews elizabethandrews left a comment

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.

FE change LGTM

@elizabethandrews elizabethandrews left a comment

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.

Oh wait I was too quick to approve. I see some of the other deprecated attributes using diagnoseDeprecatedAttribute. I don't recall what this does. Can you clarify why this wasn't used?

@softwaresat

Copy link
Copy Markdown
Contributor Author

Oh wait I was too quick to approve. I see some of the other deprecated attributes using diagnoseDeprecatedAttribute. I don't recall what this does. Can you clarify why this wasn't used?

diagnoseDeprecatedAttribute was not used because intel::named_sub_group_size has no replacement attribute spelling. The proposed replacement is a kernel property, so I emit only warn_attribute_spelling_deprecated and intentionally provide no fix-it.

@softwaresat
softwaresat force-pushed the remove-oneapi-sub-group branch from 2980106 to 4adb3b7 Compare August 30, 2026 04:10

@tahonermann tahonermann left a comment

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.

I think the diagnostic could be more helpful; see inline comments. The PR title and summary could use updates to clarify what change is being made; neither of those currently indicate this change has anything to do with attributes.

Comment on lines +133 to +134
Diag(A.getLoc(), diag::warn_attribute_spelling_deprecated)
<< "'" + A.getNormalizedFullName() + "'";

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.

The warn_attribute_spelling_deprecated diagnostic doesn't seem appropriate here. That diagnostic is currently used when only some spellings of the attribute are deprecated. For example, all [[cl::XXX]] attribute spellings have been deprecated in favor of their equivalent [[sycl::XXX]] attribute spellings. The actual message produced for the warn_attribute_spelling_deprecated is reasonable for this new deprecation, but I think it would be better to tell users what they should use instead. See the warn_atl_uuid_deprecated, warn_vector_mode_deprecated, and warn_amdgpu_num_reg_attr_deprecated diagnostics as good examples that provide the user more direction.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

Remove deprecated oneapi extensions

5 participants