Skip to content

Commit 2667e3e

Browse files
authored
[SYCL] Implement SYCL 2020 spec functionality: no propagation from function to the caller
In SYCL 1.2.1 spec, the attributes get propagated from device functions to a kernel. The SYCL 2020 requirement mandating the avoidance of the propagation of all kernel attributes to the caller when used on a function. Attributes that should not be propagated from device functions to a kernel to match with new SYCL 2020 spec. 1. scheduler_target_fmax_mhz 2. kernel_args_restrict 3. no_global_work_offset 4. max-work-group-size 5. max-global-work-dim 6. num-simd-work-items 7. reqd-sub-group-size 8. reqd-work-group-size 9. named_sub_group_size 10. sycl_explicit_simd This patch i. keeps the SYCL 1.2.1 spec functionality and propagates the attributes with the older SYCL mode(-sycl-std=2017) iii. adds or updates tests to validate the propagating behavior with SYCL 1.2.1 and SYCL 2020 specs.
1 parent a8fe4a5 commit 2667e3e

29 files changed

+1023
-98
lines changed

clang/include/clang/Basic/AttrDocs.td

Lines changed: 36 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -415,6 +415,10 @@ def SYCLSimdDocs : Documentation {
415415
The compiler may decide to compile such functions using different optimization
416416
and code generation pipeline. Also, this attribute is used to distinguish
417417
ESIMD private globals from regular SYCL global variables.
418+
419+
In SYCL 1.2.1 mode, the ``intel::sycl_explicit_simd`` attribute is propagated
420+
from the function it is applied to onto the kernel which calls the function.
421+
In SYCL 2020 mode, the attribute is not propagated to the kernel.
418422
}];
419423
}
420424

@@ -2443,8 +2447,9 @@ lambda capture, or function object member, of the callable to which the
24432447
attribute was applied. This effect is equivalent to annotating restrict on
24442448
**all** kernel pointer arguments in an OpenCL or SPIR-V kernel.
24452449

2446-
If ``intel::kernel_args_restrict`` is applied to a function called from a device
2447-
kernel, the attribute is not ignored and it is propagated to the kernel.
2450+
In SYCL 1.2.1 mode, the ``intel::kernel_args_restrict`` attribute is propagated
2451+
from the function it is applied to onto the kernel which calls the function.
2452+
In SYCL 2020 mode, the attribute is not propagated to the kernel.
24482453

24492454
The attribute forms an unchecked assertion, in that implementations
24502455
do not need to check/confirm the pre-condition in any way. If a user applies
@@ -2482,8 +2487,10 @@ def SYCLIntelNumSimdWorkItemsAttrDocs : Documentation {
24822487
let Content = [{
24832488
Applies to a device function/lambda function. Indicates the number of work
24842489
items that should be processed in parallel. Valid values are positive integers.
2485-
If ``intel::num_simd_work_items`` is applied to a function called from a
2486-
device kernel, the attribute is not ignored and it is propagated to the kernel.
2490+
2491+
In SYCL 1.2.1 mode, the ``intel::num_simd_work_items`` attribute is propagated
2492+
from the function it is applied to onto the kernel which calls the function.
2493+
In SYCL 2020 mode, the attribute is not propagated to the kernel.
24872494

24882495
.. code-block:: c++
24892496

@@ -2656,6 +2663,11 @@ allows the Y and Z arguments to be optional. If not provided by the user, the
26562663
value of Y and Z defaults to 1. See section 5.8.1 Kernel Attributes for more
26572664
details.
26582665

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.
2670+
26592671
.. code-block:: c++
26602672

26612673
[[sycl::reqd_work_group_size(4, 4, 4)]] void foo() {}
@@ -2800,8 +2812,10 @@ Applies to a device function/lambda function. Indicates the maximum dimensions
28002812
of a work group. Values must be positive integers. This is similar to
28012813
reqd_work_group_size, but allows work groups that are smaller or equal to the
28022814
specified sizes.
2803-
If ``intel::max_work_group_size`` is applied to a function called from a
2804-
device kernel, the attribute is not ignored and it is propagated to the kernel.
2815+
2816+
In SYCL 1.2.1 mode, the ``intel::max_work_group_size`` attribute is propagated
2817+
from the function it is applied to onto the kernel which calls the function.
2818+
In SYCL 2020 mode, the attribute is not propagated to the kernel.
28052819

28062820
.. code-block:: c++
28072821

@@ -2832,8 +2846,10 @@ Applies to a device function/lambda function or function call operator (of a
28322846
function object). Indicates the largest valid global work dimension that will be
28332847
accepted when running the kernel on a device. Valid values are integers in a
28342848
range of [0, 3].
2835-
If ``intel::max_global_work_dim`` is applied to a function called from a
2836-
device kernel, the attribute is not ignored and it is propagated to the kernel.
2849+
2850+
In SYCL 1.2.1 mode, the ``intel::max_global_work_dim`` attribute is propagated
2851+
from the function it is applied to onto the kernel which calls the function.
2852+
In SYCL 2020 mode, the attribute is not propagated to the kernel.
28372853

28382854
.. code-block:: c++
28392855

@@ -2890,6 +2906,10 @@ device operation, guiding the FPGA backend to insert the appropriate number of
28902906
registers to break-up the combinational logic circuit, and thereby controlling
28912907
the length of the longest combinational path.
28922908

2909+
In SYCL 1.2.1 mode, the ``intel::scheduler_target_fmax_mhz`` attribute is
2910+
propagated from the function it is applied to onto the kernel which calls the
2911+
function. In SYCL 2020 mode, the attribute is not propagated to the kernel.
2912+
28932913
.. code-block:: c++
28942914

28952915
[[intel::scheduler_target_fmax_mhz(4)]] void foo() {}
@@ -2920,6 +2940,10 @@ function object). If 1, compiler doesn't use the global work offset values for
29202940
the device function. Valid values are 0 and 1. If used without argument, value
29212941
of 1 is set implicitly.
29222942

2943+
In SYCL 1.2.1 mode, the ``intel::no_global_work_offset`` attribute is
2944+
propagated from the function it is applied to onto the kernel which calls the
2945+
function. In SYCL 2020 mode, the attribute is not propagated to the kernel.
2946+
29232947
.. code-block:: c++
29242948

29252949
[[intel::no_global_work_offset]]
@@ -4607,6 +4631,10 @@ the ``[[intel::named_sub_group_size(NAME)]]`` documentation for clarification.
46074631
This attribute is mutually exclusive with ``[[intel::named_sub_group_size(NAME)]]``
46084632
and ``[[intel::sycl_explicit_simd]]``.
46094633

4634+
In SYCL 1.2.1 mode, the ``intel::reqd_sub_group_size`` attribute is propagated
4635+
from the function it is applied to onto the kernel which calls the function.
4636+
In SYCL 2020 mode, the attribute is not propagated to the kernel.
4637+
46104638
In addition to device functions, the required sub-group size attribute may also
46114639
be specified in the definition of a named functor object and lambda functions,
46124640
as in the examples below:

clang/lib/Sema/SemaSYCL.cpp

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

561-
llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
562-
// FIXME: Make this list self-adapt as new SYCL attributes are added.
563-
return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
564-
ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
565-
SYCLIntelNumSimdWorkItemsAttr,
566-
SYCLIntelSchedulerTargetFmaxMhzAttr,
567-
SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
568-
SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
569-
});
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 || S.getASTContext().getLangOpts().getSYCLVersion() <
565+
LangOptions::SYCL_2020) {
566+
llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
567+
// FIXME: Make this list self-adapt as new SYCL attributes are added.
568+
return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
569+
ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
570+
SYCLIntelNumSimdWorkItemsAttr,
571+
SYCLIntelSchedulerTargetFmaxMhzAttr,
572+
SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
573+
SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
574+
});
575+
}
570576

571577
// Attributes that should not be propagated from device functions to a kernel.
572578
if (DirectlyCalled) {

0 commit comments

Comments
 (0)