Skip to content

[SYCL] PR 3 - Remove FPGA attributes from SYCL FE #21735

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
KornevNikita merged 4 commits into from
Apr 15, 2026
Merged

[SYCL] PR 3 - Remove FPGA attributes from SYCL FE #21735

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b1ef304
Resolve merge conflicts
premanandrao Apr 14, 2026
1bfa7e0
Resolve merge conflicts
premanandrao Apr 14, 2026
3f1499d
Resolve merge conflicts
premanandrao Apr 14, 2026
0c7ab82
Apply clang-format counter-intuitive suggestion
premanandrao Apr 14, 2026
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
120 changes: 0 additions & 120 deletions clang/include/clang/Basic/Attr.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2915,73 +2915,6 @@ def Mode : Attr {
let PragmaAttributeSupport = 0;
}

def SYCLIntelIVDep : StmtAttr {
Copy link
Copy Markdown
Contributor

@Fznamznon Fznamznon Apr 13, 2026

Choose a reason for hiding this comment

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

I was curious about what ivdep is and google gave me this https://www.intel.com/content/www/us/en/docs/oneapi-fpga-add-on/developer-guide/2025-0/ivdep-attribute.html . I wonder if need to ask someone to remove this from somewhere too

let Spellings = [CXX11<"intel", "ivdep">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
ErrorDiag, "'for', 'while', and 'do' statements">;
let Args = [
ExprArgument<"SafelenExpr", /*opt*/1>, ExprArgument<"ArrayExpr", /*opt*/1>,
UnsignedArgument<"SafelenValue", /*opt*/1>
];
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let AdditionalMembers = [{
bool isDependent() const {
return (getSafelenExpr() &&
getSafelenExpr()->isInstantiationDependent()) ||
(getArrayExpr() && getArrayExpr()->isInstantiationDependent());
}

const ValueDecl *getArrayDecl() const {
const Expr* E = getArrayExpr();
if (!E) return nullptr;

if (const auto *DRE = dyn_cast<DeclRefExpr>(E))
return cast<ValueDecl>(DRE->getDecl()->getCanonicalDecl());

return cast<ValueDecl>(
cast<MemberExpr>(E)->getMemberDecl()->getCanonicalDecl());
}

bool isInf() const {
return !getSafelenExpr();
}

static bool SafelenCompare(const SYCLIntelIVDepAttr *LHS,
const SYCLIntelIVDepAttr *RHS) {
// INF < INF is false, !INF < INF is true.
if (!RHS->getSafelenExpr())
return false;
if (!LHS->getSafelenExpr())
return true;
return LHS->getSafelenValue() > RHS->getSafelenValue();
}
}];
let Documentation = [SYCLIntelIVDepAttrDocs];
Copy link
Copy Markdown
Contributor

@Fznamznon Fznamznon Apr 13, 2026

Choose a reason for hiding this comment

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

Wait, should we also remove all the docs?

Copy link
Copy Markdown
Contributor Author

@premanandrao premanandrao Apr 13, 2026

Choose a reason for hiding this comment

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

Thanks, I removed from AttrDocs.td. I will go back and see if I missed removing any from the previous two PRs; and if it is from the merged one, will create a new PR for that.

}

def SYCLIntelInitiationInterval : DeclOrStmtAttr {
let Spellings = [CXX11<"intel", "ii">,
CXX11<"intel", "initiation_interval">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt, Function],
ErrorDiag,
"'for', 'while', 'do' statements, and functions">;
let Args = [ExprArgument<"NExpr">];
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Documentation = [SYCLIntelInitiationIntervalAttrDocs];
let SupportsNonconformingLambdaSyntax = 1;
}

def SYCLIntelMaxConcurrency : DeclOrStmtAttr {
let Spellings = [CXX11<"intel", "max_concurrency">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt, Function],
ErrorDiag,
"'for', 'while', 'do' statements, and functions">;
let Args = [ExprArgument<"NExpr">];
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Documentation = [SYCLIntelMaxConcurrencyAttrDocs];
let SupportsNonconformingLambdaSyntax = 1;
}

def SYCLIntelLoopCoalesce : StmtAttr {
let Spellings = [CXX11<"intel", "loop_coalesce">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
Expand All @@ -2991,26 +2924,6 @@ def SYCLIntelLoopCoalesce : StmtAttr {
let Documentation = [SYCLIntelLoopCoalesceAttrDocs];
}

def SYCLIntelDisableLoopPipelining : DeclOrStmtAttr {
let Spellings = [CXX11<"intel", "disable_loop_pipelining">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt, Function],
ErrorDiag,
"'for', 'while', 'do' statements, and functions">;
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Documentation = [SYCLIntelDisableLoopPipeliningAttrDocs];
let SupportsNonconformingLambdaSyntax = 1;
let SimpleHandler = 1;
}
def : MutualExclusions<[SYCLIntelInitiationInterval,
SYCLIntelDisableLoopPipelining]>;
def : MutualExclusions<[SYCLIntelIVDep,
SYCLIntelDisableLoopPipelining]>;
def : MutualExclusions<[SYCLIntelMaxConcurrency,
SYCLIntelDisableLoopPipelining]>;

def : MutualExclusions<[SYCLIntelMaxConcurrency,
SYCLIntelDisableLoopPipelining]>;

def SYCLIntelMaxInterleaving : StmtAttr {
let Spellings = [CXX11<"intel", "max_interleaving">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
Expand All @@ -3020,39 +2933,6 @@ def SYCLIntelMaxInterleaving : StmtAttr {
let Documentation = [SYCLIntelMaxInterleavingAttrDocs];
}

def SYCLIntelSpeculatedIterations : StmtAttr {
let Spellings = [CXX11<"intel", "speculated_iterations">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
ErrorDiag, "'for', 'while', and 'do' statements">;
let Args = [ExprArgument<"NExpr">];
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Documentation = [SYCLIntelSpeculatedIterationsAttrDocs];
}
def : MutualExclusions<[SYCLIntelDisableLoopPipelining,
SYCLIntelSpeculatedIterations]>;

def SYCLIntelMaxReinvocationDelay : StmtAttr {
let Spellings = [CXX11<"intel", "max_reinvocation_delay">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
ErrorDiag, "'for', 'while', and 'do' statements">;
let Args = [ExprArgument<"NExpr">];
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Documentation = [SYCLIntelMaxReinvocationDelayAttrDocs];
}
def : MutualExclusions<[SYCLIntelDisableLoopPipelining,
SYCLIntelMaxReinvocationDelay]>;

def SYCLIntelEnableLoopPipelining : StmtAttr {
let Spellings = [CXX11<"intel", "enable_loop_pipelining">];
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
ErrorDiag, "'for', 'while', and 'do' statements">;
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Documentation = [SYCLIntelEnableLoopPipeliningAttrDocs];
}

def : MutualExclusions<[SYCLIntelDisableLoopPipelining,
SYCLIntelEnableLoopPipelining]>;

def SYCLIntelDoublePump : Attr {
let Spellings = [CXX11<"intel", "doublepump">];
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
Expand Down
210 changes: 0 additions & 210 deletions clang/include/clang/Basic/AttrDocs.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4013,108 +4013,6 @@ with which a pipe interfaces. The id argument is the name of the I/O interface.
}];
}

def SYCLIntelIVDepAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::ivdep";
let Content = [{
This attribute applies to a loop. If no additional arguments are provided, it
indicates that the backend may assume that the loop carries no dependences.
If a safelen is provided then the backend may assume that all loop-carried
dependences have a dependence distance of at least that length.
If an array variable is provided then the backend may assume that there are
no loop-carried dependences for this particular array within this loop.
If both an array variable and a safelen are specified, the backend may assume
that all loop-carried dependences for this particular array have a dependence
distance of at least that length.
Multiple ivdep attributes can be applied to the same loop, but only the greatest
possible safelen will be chosen for each array, including the case where safelen
is unspecified and is therefore considered infinite.
In case of ivdep being applied both w/o an array variable and for a particular
array, the array variables that were not designated a separate ivdep will receive
the no-array ivdep's safelen, with the correspondent treatment by the backend.
The ``[[intel::ivdep]]`` attribute applies a safelen value of 0 or 1 that has
no effect on the loop and emits a suppressible warning when safelen value of 0
or 1 is used.

.. code-block:: c++

void foo() {
int a[10];
[[intel::ivdep]] for (int i = 0; i != 10; ++i) { }
[[intel::ivdep(2)]] for (int i = 0; i != 10; ++i) { }
[[intel::ivdep(a)]] for (int i = 0; i != 10; ++i) { }
[[intel::ivdep(a, 2)]] for (int i = 0; i != 10; ++i) { }
}

template<int N>
void bar() {
[[intel::ivdep(N)]] for(;;) { }
}

}];
}

def SYCLIntelInitiationIntervalAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::initiation_interval";
let Content = [{
This attribute applies to a loop or a function. Indicates that the loop or
function should be pipelined with an initiation interval of N. N must be a
positive integer. Cannot be applied multiple times to the same loop or function.
Cannot be used on the same loop or function in conjunction with
disable_loop_pipelining.

The ``[[intel::ii]]`` attribute spelling is a deprecated synonym for
``[[intel::initiation_interval]]`` and will be removed in the future.

.. code-block:: c++

void foo() {
int var = 0;
[[intel::initiation_interval(4)]] for (int i = 0; i < 10; ++i) var++;
}

[[intel::initiation_interval(4)]] void foo1 { }

template<int N>
void bar() {
[[intel::initiation_interval(N)]] for(;;) { }
}

template<int N>
[[intel::initiation_interval(N)]] void bar1 { }

}];
}

def SYCLIntelMaxConcurrencyAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::max_concurrency";
let Content = [{
This attribute applies to a loop or a function. It indicates that the
loop/function should allow no more than N threads or iterations to execute it
simultaneously. N must be a non negative integer. '0' indicates the
max_concurrency case to be unbounded. Cannot be applied multiple times to the
same loop or function, or in conjunction with ``disable_loop_pipelining``.

.. code-block:: c++

void foo() {
int a[10];
[[intel::max_concurrency(2)]] for (int i = 0; i != 10; ++i) a[i] = 0;
}

[[intel::max_concurrency(2)]] void foo1 { }
template<int N>
void bar() {
[[intel::max_concurrency(N)]] for(;;) { }
}
template<int N>
[[intel::max_concurrency(N)]] void bar1() { }

}];
}

def SYCLIntelLoopCoalesceAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::loop_coalesce";
Expand Down Expand Up @@ -4154,28 +4052,6 @@ of the nested loop levels should be coalesced.
}];
}

def SYCLIntelDisableLoopPipeliningAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::disable_loop_pipelining";
let Content = [{
This attribute applies to a loop or a function. Takes no arguments and
indicates whether the kernel should be pipelined or not. Cannot be used on the
same loop or function, or in conjunction with ``speculated_iterations``,
``max_concurrency``, ``initiation_interval``, ``ivdep``,
``max_reinvocation_delay`` or ``enable_loop_pipelining`` attribute.

.. code-block:: c++

void foo() {
int var = 0;
[[intel::disable_loop_pipelining]] for (int i = 0; i < 10; ++i) var++;
}

[[intel::disable_loop_pipelining]] void foo1() { }

}];
}

def SYCLIntelMaxInterleavingAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::max_interleaving";
Expand All @@ -4201,92 +4077,6 @@ Parameter N is mandatory, and may either be 0, or 1.
}];
}

def SYCLIntelSpeculatedIterationsAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::speculated_iterations";
let Content = [{
This attribute applies to a loop. Specifies the number of concurrent speculated
iterations that will be in flight for a loop invocation (i.e. the exit
condition for these iterations will not have been evaluated yet).
Parameter N is mandatory, and may either be 0, or a positive integer. Cannot be
used on the same loop in conjunction with disable_loop_pipelining.

.. code-block:: c++

void foo() {
int var = 0;
[[intel::speculated_iterations(4)]] for (int i = 0; i < 10; ++i) var++;
}

template<int N>
void bar() {
[[intel::speculated_iterations(N)]] for(;;) { }
}

}];
}

def SYCLIntelMaxReinvocationDelayAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::max_reinvocation_delay";
let Content = [{
This attribute applies to a loop. Specifies the maximum number of cycles allowed
on the delay between the launch of the last iteration of a loop invocation and
the launch of the first iteration of a new loop invocation. Parameter N is
mandatory, and is a positive integer. Cannot be used on the same loop in
conjunction with disable_loop_pipelining.

.. code-block:: c++

void foo() {
int var = 0;
[[intel::max_reinvocation_delay(1)]]
for (int i = 0; sycl::log10((float)(x)) < 10; i++) var++;
}

template<int N>
void bar() {
[[intel::max_reinvocation_delay(N)]] for(;;) { }
}
}];
}

def SYCLIntelEnableLoopPipeliningAttrDocs : Documentation {
let Category = DocCatVariable;
let Heading = "intel::enable_loop_pipelining";
let Content = [{
The ``enable_loop_pipelining`` attribute applies to a loop in SYCL device code.
Takes no arguments and enables pipelining of the loop. This attribute is useful
in the low-area flow, in which all loops are unpipelined by default. Cannot be
used on the same loop in conjunction with ``disable_loop_pipelining`` attribute.

.. code-block:: c++

void bar() {
int a[10];
[[[intel::enable_loop_pipelining]] for (int i = 0; i != 10; ++i) a[i] = 0;
}

void Array(int *array, size_t n) {
[[intel::enable_loop_pipelining]] for (int i = 0; i < n; ++i) array[i] = 0;
}

void count () {
int a1[10], i = 0;
[[intel::enable_loop_pipelining]] while (i < 10) {
a1[i++] = 3;
}
}

void check() {
int a = 10;
[[intel::enable_loop_pipelining]] do {
a = a + 1;
} while (a < 20);
}

}];
}

def SYCLAddIRAttributesFunctionDocs : Documentation {
let Category = DocCatFunction;
Expand Down
12 changes: 0 additions & 12 deletions clang/include/clang/Sema/Sema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11405,27 +11405,15 @@ class Sema final : public SemaBase {
const IdentifierInfo *AttrName,
SourceRange Range);

SYCLIntelIVDepAttr *BuildSYCLIntelIVDepAttr(const AttributeCommonInfo &CI,
Expr *Expr1, Expr *Expr2);
LoopUnrollHintAttr *BuildLoopUnrollHintAttr(const AttributeCommonInfo &A,
Expr *E);
OpenCLUnrollHintAttr *
BuildOpenCLLoopUnrollHintAttr(const AttributeCommonInfo &A, Expr *E);

SYCLIntelInitiationIntervalAttr *
BuildSYCLIntelInitiationIntervalAttr(const AttributeCommonInfo &CI, Expr *E);
SYCLIntelMaxConcurrencyAttr *
BuildSYCLIntelMaxConcurrencyAttr(const AttributeCommonInfo &CI, Expr *E);
SYCLIntelMaxInterleavingAttr *
BuildSYCLIntelMaxInterleavingAttr(const AttributeCommonInfo &CI, Expr *E);
SYCLIntelSpeculatedIterationsAttr *
BuildSYCLIntelSpeculatedIterationsAttr(const AttributeCommonInfo &CI,
Expr *E);
SYCLIntelLoopCoalesceAttr *
BuildSYCLIntelLoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E);
SYCLIntelMaxReinvocationDelayAttr *
BuildSYCLIntelMaxReinvocationDelayAttr(const AttributeCommonInfo &CI,
Expr *E);

///@}

Expand Down
Loading
Loading