Skip to content

Commit fcc5de0

Browse files
committed
address @aaron's review comments
Signed-off-by: Soumi Manna <[email protected]>
1 parent 79f9bf1 commit fcc5de0

File tree

2 files changed

+9
-21
lines changed

2 files changed

+9
-21
lines changed

clang/include/clang/Basic/AttrDocs.td

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2663,10 +2663,10 @@ allows the Y and Z arguments to be optional. If not provided by the user, the
26632663
value of Y and Z defaults to 1. See section 5.8.1 Kernel Attributes for more
26642664
details.
26652665

2666-
In SYCL 1.2.1 mode, the ``intel::reqd_work_group_size`` or
2667-
``cl::reqd_work_group_size`` or ``sycl::reqd_work_group_size`` attribute is
2668-
propagated from the function it is applied to onto the kernel which calls the
2669-
function. In SYCL 2020 mode, the attribute is not propagated to the kernel.
2666+
In SYCL 1.2.1 mode, the ``intel::reqd_work_group_size``,
2667+
``cl::reqd_work_group_size``, and ``sycl::reqd_work_group_size`` attributes are
2668+
propagated from the function they are applied to onto the kernel which calls the
2669+
function. In SYCL 2020 mode, the attributes are not propagated to the kernel.
26702670

26712671
.. code-block:: c++
26722672

clang/lib/Sema/SemaSYCL.cpp

Lines changed: 5 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -558,10 +558,11 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
558558
if (!FD->hasAttrs())
559559
return;
560560

561-
// Attributes that should be propagated from device functions to a kernel
562-
// in SYCL 1.2.1.
563-
if (S.getASTContext().getLangOpts().getSYCLVersion() <
564-
LangOptions::SYCL_2020) {
561+
// In SYCL 1.2.1 mode, the attributes are propagated from the function they
562+
// are applied to onto the kernel which calls the function.
563+
// In SYCL 2020 mode, the attributes are not propagated to the kernel.
564+
if (DirectlyCalled ||
565+
S.getASTContext().getLangOpts().getSYCLVersion() < LangOptions::SYCL_2020) {
565566
llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
566567
// FIXME: Make this list self-adapt as new SYCL attributes are added.
567568
return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
@@ -571,19 +572,6 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
571572
SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
572573
SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
573574
});
574-
} else {
575-
// Attributes that should not be propagated from device functions to a
576-
// kernel in SYCL 2020.
577-
if (DirectlyCalled) {
578-
llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
579-
return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
580-
ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
581-
SYCLIntelNumSimdWorkItemsAttr,
582-
SYCLIntelSchedulerTargetFmaxMhzAttr,
583-
SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
584-
SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
585-
});
586-
}
587575
}
588576

589577
// Attributes that should not be propagated from device functions to a kernel.

0 commit comments

Comments
 (0)