-
Notifications
You must be signed in to change notification settings - Fork 18.6k
Expand file tree
/
Copy pathAttrDocs.td
More file actions
10652 lines (8889 loc) · 403 KB
/
Copy pathAttrDocs.td
File metadata and controls
10652 lines (8889 loc) · 403 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
687
688
689
690
691
692
693
694
695
696
697
698
699
700
701
702
703
704
705
706
707
708
709
710
711
712
713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
760
761
762
763
764
765
766
767
768
769
770
771
772
773
774
775
776
777
778
779
780
781
782
783
784
785
786
787
788
789
790
791
792
793
794
795
796
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
893
894
895
896
897
898
899
900
901
902
903
904
905
906
907
908
909
910
911
912
913
914
915
916
917
918
919
920
921
922
923
924
925
926
927
928
929
930
931
932
933
934
935
936
937
938
939
940
941
942
943
944
945
946
947
948
949
950
951
952
953
954
955
956
957
958
959
960
961
962
963
964
965
966
967
968
969
970
971
972
973
974
975
976
977
978
979
980
981
982
983
984
985
986
987
988
989
990
991
992
993
994
995
996
997
998
999
1000
//==--- AttrDocs.td - Attribute documentation ----------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===---------------------------------------------------------------------===//
// To test that the documentation builds cleanly, you must run clang-tblgen to
// convert the .td file into a .md file, and then run sphinx to convert the
// .md file into an HTML file. After completing testing, you should revert the
// generated .md file so that the modified version does not get checked in to
// version control.
//
// To run clang-tblgen to generate the .md file:
// clang-tblgen -gen-attr-docs -I <root>/clang/include
// <root>/clang/include/clang/Basic/Attr.td -o
// <root>/clang/docs/AttributeReference.md
//
// To run sphinx to generate the .html files (note that sphinx-build must be
// available on the PATH):
// Windows (from within the clang\docs directory):
// make.bat html
// Non-Windows (from within the clang/docs directory):
// sphinx-build -b html . _build/html
def GlobalDocumentation {
code Intro =[{% -------------------------------------------------------------------
% NOTE: This file is automatically generated by running clang-tblgen
% -gen-attr-docs. Do not edit this file by hand!!
% -------------------------------------------------------------------
# Attributes in Clang
```{program} clang
```
## Introduction
This page lists the attributes currently supported by Clang.
}];
}
def SectionDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The `section` attribute allows you to specify a specific section a
global variable or function should be in after translation.
}];
let Heading = "section, __declspec(allocate)";
}
def CodeModelDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The `model` attribute allows overriding the translation unit's
code model (specified by `-mcmodel`) for a specific global variable.
On LoongArch, allowed values are "normal", "medium", "extreme".
On x86-64, allowed values are `"small"` and `"large"`. `"small"` is
roughly equivalent to `-mcmodel=small`, meaning the global is considered
"small" placed closer to the `.text` section relative to "large" globals, and
to prefer using 32-bit relocations to access the global. `"large"` is roughly
equivalent to `-mcmodel=large`, meaning the global is considered "large" and
placed further from the `.text` section relative to "small" globals, and
64-bit relocations must be used to access the global.
}];
let Heading = "model";
}
def UsedDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
This attribute, when attached to a function or variable definition, indicates
that there may be references to the entity which are not apparent in the source
code. For example, it may be referenced from inline `asm`, or it may be
found through a dynamic symbol or section lookup.
The compiler must emit the definition even if it appears to be unused, and it
must not apply optimizations which depend on fully understanding how the entity
is used.
Whether this attribute has any effect on the linker depends on the target and
the linker. Most linkers support the feature of section garbage collection
(`--gc-sections`), also known as "dead stripping" (`ld64 -dead_strip`) or
discarding unreferenced sections (`link.exe /OPT:REF`). On COFF and Mach-O
targets (Windows and Apple platforms), the `used` attribute prevents symbols
from being removed by linker section GC. On ELF targets, it has no effect on its
own, and the linker may remove the definition if it is not otherwise referenced.
This linker GC can be avoided by also adding the `retain` attribute. Note
that `retain` requires special support from the linker; see that attribute's
documentation for further information.
}];
}
def RetainDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
This attribute, when attached to a function or variable definition, prevents
section garbage collection in the linker. It does not prevent other discard
mechanisms, such as archive member selection, and COMDAT group resolution.
If the compiler does not emit the definition, e.g. because it was not used in
the translation unit or the compiler was able to eliminate all of the uses,
this attribute has no effect. This attribute is typically combined with the
`used` attribute to force the definition to be emitted and preserved into the
final linked image.
This attribute is only necessary on ELF targets; other targets prevent section
garbage collection by the linker when using the `used` attribute alone.
Using the attributes together should result in consistent behavior across
targets.
This attribute requires the linker to support the `SHF_GNU_RETAIN` extension.
This support is available in GNU `ld` and `gold` as of binutils 2.36, as
well as in `ld.lld` 13.
}];
}
def InitPriorityDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
In C++, the order in which global variables are initialized across translation
units is unspecified, unlike the ordering within a single translation unit. The
`init_priority` attribute allows you to specify a relative ordering for the
initialization of objects declared at namespace scope in C++ within a single
linked image on supported platforms. The priority is given as an integer constant
expression between 101 and 65535 (inclusive). Priorities outside of that range are
reserved for use by the implementation. A lower value indicates a higher priority
of initialization. Note that only the relative ordering of values is important.
For example:
```c++
struct SomeType { SomeType(); };
__attribute__((init_priority(200))) SomeType Obj1;
__attribute__((init_priority(101))) SomeType Obj2;
```
`Obj2` will be initialized *before* `Obj1` despite the usual order of
initialization being the opposite.
Note that this attribute does not control the initialization order of objects
across final linked image boundaries like shared objects and executables.
On Windows, `init_seg(compiler)` is represented with a priority of 200 and
`init_seg(library)` is represented with a priority of 400. `init_seg(user)`
uses the default 65535 priority.
On MachO platforms, this attribute also does not control the order of initialization
across translation units, where it only affects the order within a single TU.
This attribute is only supported for C++ and Objective-C++ and is ignored in
other language modes.
}];
}
def InitSegDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The attribute applied by `pragma init_seg()` controls the section into
which global initialization function pointers are emitted. It is only
available with `-fms-extensions`. Typically, this function pointer is
emitted into `.CRT$XCU` on Windows. The user can change the order of
initialization by using a different section name with the same
`.CRT$XC` prefix and a suffix that sorts lexicographically before or
after the standard `.CRT$XCU` sections. See the [init_seg][init_seg]
documentation on MSDN for more information.
[init_seg]: http://msdn.microsoft.com/en-us/library/7977wcck(v=vs.110).aspx
}];
}
def TLSModelDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The `tls_model` attribute allows you to specify which thread-local storage
model to use. It accepts the following strings:
- global-dynamic
- local-dynamic
- initial-exec
- local-exec
TLS models are mutually exclusive.
}];
}
def AddressSpaceDocs : Documentation {
let Category = DocCatType;
let Heading = "address_space";
let Label = "langext-address_space_documentation";
let Content = [{
:::{Note}
This attribute is mainly intended to be used by target headers
provided by the toolchain. End users should prefer the documented, named
address space annotations for their platform, such as the
[OpenCL address spaces], `__global__`, `__local__`, or something else.
:::
The `address_space` attribute functions as a type qualifier that allows the
programmer to specify the address space for a pointer or reference type.
Qualified pointer types are considered distinct types for the purposes of
overload resolution. The attribute takes a single, non-negative integer
constant expression identifying the address space. For example:
```c
int * __attribute__((address_space(1))) ptr;
void foo(__attribute__((address_space(2))) float *buf);
```
Only one address space qualifier may be applied to a given pointer or reference
type. Where address spaces are allowed (e.g., variables, parameters, return
types) and what values are valid depends on the target and language mode.
The meaning of each value is defined by the target; multiple address spaces are
used in environments such as OpenCL, CUDA, HIP, and other GPU programming
models to distinguish global, local, constant, and private memory. See for
example the address spaces defined in the [NVPTX Usage Guide][nvptx usage guide] and the
[AMDGPU Usage Guide][amdgpu usage guide].
Address spaces may partially overlap or be entirely distinct. The compiler may
reject attempts to convert between distinct, incompatible address spaces.
Pointer width may vary between different address spaces, so some explicit casts
may truncate.
For more information, refer to [ISO TR18037][iso tr18037], which covers embedded C language
extensions. Section 5 covers named address spaces.
[amdgpu usage guide]: https://llvm.org/docs/AMDGPUUsage.html#address-spaces
[iso tr18037]: https://standards.iso.org/ittf/PubliclyAvailableStandards/c051126_ISO_IEC_TR_18037_2008.zip
[nvptx usage guide]: https://llvm.org/docs/NVPTXUsage.html#address-spaces
}];
}
def DLLExportDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The `__declspec(dllexport)` attribute declares a variable, function, or
Objective-C interface to be exported from the module. It is available under the
`-fdeclspec` flag for compatibility with various compilers. The primary use
is for COFF object files which explicitly specify what interfaces are available
for external use. See the [dllexport][dllexport] documentation on MSDN for more
information.
[dllexport]: https://msdn.microsoft.com/en-us/library/3y1sfaz2.aspx
}];
}
def DLLImportDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The `__declspec(dllimport)` attribute declares a variable, function, or
Objective-C interface to be imported from an external module. It is available
under the `-fdeclspec` flag for compatibility with various compilers. The
primary use is for COFF object files which explicitly specify what interfaces
are imported from external modules. See the [dllimport][dllimport] documentation on MSDN
for more information.
Note that a dllimport function may still be inlined, if its definition is
available and it doesn't reference any non-dllimport functions or global
variables.
[dllimport]: https://msdn.microsoft.com/en-us/library/3y1sfaz2.aspx
}];
}
def ThreadDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The `__declspec(thread)` attribute declares a variable with thread local
storage. It is available under the `-fms-extensions` flag for MSVC
compatibility. See the documentation for [\_\_declspec(thread)][__declspec(thread)] on MSDN.
In Clang, `__declspec(thread)` is generally equivalent in functionality to the
GNU `__thread` keyword. The variable must not have a destructor and must have
a constant initializer, if any. The attribute only applies to variables
declared with static storage duration, such as globals, class static data
members, and static locals.
[__declspec(thread)]: http://msdn.microsoft.com/en-us/library/9w1sdazb.aspx
}];
}
def NoEscapeDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
`noescape` placed on a function parameter of a pointer type is used to inform
the compiler that the pointer cannot escape: that is, no reference to the object
the pointer points to that is derived from the parameter value will survive
after the function returns. Users are responsible for making sure parameters
annotated with `noescape` do not actually escape. The optimizer may make
assumptions based on the fact that it knows that a call to the function does
not escape a certain parameter, so incorrectly annotating a parameter with
`noescape` leads to undefined behavior. The callee is also not allowed to
deallocate memory through a `noescape` parameter: the optimizer does not make
assumptions based on this information at the moment, but may do so in the
future. Some cases of invalid uses of `noescape` can be found with
{ref}`-Wlifetime-safety-noescape <Wlifetime-safety-noescape>`.
For example:
```c
int *gp;
void nonescapingFunc(__attribute__((noescape)) int *p) {
*p += 100; // OK.
}
void escapingFunc(__attribute__((noescape)) int *p) {
gp = p; // Not OK.
}
void freeingFunc(__attribute__((noescape)) int *p) {
free(p); // Not OK.
}
```
Since `noescape` is a parameter attribute and not a type attribute, it only
applies to the outermost pointer level, regardless of where in the parameter
declaration you place it:
```c
int **gp;
void nestingEscapes(__attribute__((noescape)) int **p) {
gp = p; // Not OK.
*gp = *p; // OK, p does not escape.
}
```
Additionally, when the parameter is a
{doc}`block pointer <BlockLanguageSpec>`, the same restriction applies to
copies of the block. For example:
```c
typedef void (^BlockTy)();
BlockTy g0, g1;
void nonescapingFunc(__attribute__((noescape)) BlockTy block) {
block(); // OK.
}
void escapingFunc(__attribute__((noescape)) BlockTy block) {
g0 = block; // Not OK.
g1 = Block_copy(block); // Not OK either.
}
```
The function *is* allowed to leak information about the memory address of the
pointer, but not any provenance of the allocation:
```c
bool isNull(__attribute__((noescape)) void *p) {
return !p; // OK.
}
uintptr_t gi;
void escapingAddress(__attribute__((noescape)) int *p) {
// OK *if and only if* gi is never casted back to a pointer.
gi = (uintptr_t)p;
}
bool usingEscapedAddress(int *p) {
return (uintptr_t)p > gi; // OK.
}
bool usingEscapedPointer(int *p) {
return p > (int*)gi; // Not OK.
}
int *gp;
void escapingEndFunc(__attribute__((noescape)) int *p, size_t len) {
gp = p + len; // Not OK.
}
```
}];
}
def WarnUnusedDocs : Documentation {
let Category = DocCatType;
let Content = [{
The `warn_unused` attribute can be placed on the declaration of a structure or union type.
When the `-Wunused-variable` diagnostic is enabled, local variables of types which have a non-trivial constructor or destructor are considered "used" by virtue of the constructor or destructor invocations involved.
Those constructor or destructor invocations are not considered a use if the type is declared with the `warn_unused` attribute.
The variable is considered used if it is named outside of its declaration.
This attribute is available in both C and C++ language modes but is primarily useful in C++ for classes which have a non-trivial constructor or destructor but act as a value type rather than an RAII type.
```c++
struct [[gnu::warn_unused]] S {
S();
~S();
};
struct T {
T();
~T();
};
int func() {
S s1; // -Wunused-variable warning
S s2; // No -Wunused-variable warning because of the member access expression below
S s3; // No -Wunused-variable warning because of the sizeof operand below
T t; // No -Wunused-variable warning
s2.~S();
return sizeof(s3);
}
```
}];
}
def MaybeUndefDocs : Documentation {
let Category = DocCatVariable;
let Content = [{
The `maybe_undef` attribute can be placed on a function parameter. It indicates
that the parameter is allowed to use undef values. It informs the compiler
to insert a freeze LLVM IR instruction on the function parameter.
Please note that this is an attribute that is used as an internal
implementation detail and not intended to be used by external users.
In languages HIP, CUDA etc., some functions have multi-threaded semantics and
it is enough for only one or some threads to provide defined arguments.
Depending on semantics, undef arguments in some threads don't produce
undefined results in the function call. Since, these functions accept undefined
arguments, `maybe_undef` attribute can be placed.
Sample usage:
```c
void maybeundeffunc(int __attribute__((maybe_undef))param);
```
}];
}
def CPUSpecificCPUDispatchDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
The `cpu_specific` and `cpu_dispatch` attributes are used to define and
resolve multiversioned functions. This form of multiversioning provides a
mechanism for declaring versions across translation units and manually
specifying the resolved function list. A specified CPU defines a set of minimum
features that are required for the function to be called. The result of this is
that future processors execute the most restrictive version of the function the
new processor can execute.
In addition, unlike the ICC implementation of this feature, the selection of the
version does not consider the manufacturer or microarchitecture of the processor.
It tests solely the list of features that are both supported by the specified
processor and present in the compiler-rt library. This can be surprising at times,
as the runtime processor may be from a completely different manufacturer, as long
as it supports the same feature set.
This can additionally be surprising, as some processors are indistringuishable from
others based on the list of testable features. When this happens, the variant
is selected in an unspecified manner.
Function versions are defined with `cpu_specific`, which takes one or more CPU
names as a parameter. For example:
```c
// Declares and defines the ivybridge version of single_cpu.
__attribute__((cpu_specific(ivybridge)))
void single_cpu(void){}
// Declares and defines the atom version of single_cpu.
__attribute__((cpu_specific(atom)))
void single_cpu(void){}
// Declares and defines both the ivybridge and atom version of multi_cpu.
__attribute__((cpu_specific(ivybridge, atom)))
void multi_cpu(void){}
```
A dispatching (or resolving) function can be declared anywhere in a project's
source code with `cpu_dispatch`. This attribute takes one or more CPU names
as a parameter (like `cpu_specific`). Functions marked with `cpu_dispatch`
are not expected to be defined, only declared. If such a marked function has a
definition, any side effects of the function are ignored; trivial function
bodies are permissible for ICC compatibility.
```c
// Creates a resolver for single_cpu above.
__attribute__((cpu_dispatch(ivybridge, atom)))
void single_cpu(void){}
// Creates a resolver for multi_cpu, but adds a 3rd version defined in another
// translation unit.
__attribute__((cpu_dispatch(ivybridge, atom, sandybridge)))
void multi_cpu(void){}
```
Note that it is possible to have a resolving function that dispatches based on
more or fewer options than are present in the program. Specifying fewer will
result in the omitted options not being considered during resolution. Specifying
a version for resolution that isn't defined in the program will result in a
linking failure.
It is also possible to specify a CPU name of `generic` which will be resolved
if the executing processor doesn't satisfy the features required in the CPU
name. The behavior of a program executing on a processor that doesn't satisfy
any option of a multiversioned function is undefined.
}];
}
def SYCLKernelDocs : Documentation {
let Category = DocCatFunction;
let Heading = "sycl_kernel";
let Content = [{
The `sycl_kernel` attribute specifies that a function template will be used
to outline device code and to generate an OpenCL kernel.
Here is a code example of the SYCL program, which demonstrates the compiler's
outlining job:
```c++
int foo(int x) { return ++x; }
using namespace cl::sycl;
queue Q;
buffer<int, 1> a(range<1>{1024});
Q.submit([&](handler& cgh) {
auto A = a.get_access<access::mode::write>(cgh);
cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
A[index] = index[0] + foo(42);
});
}
```
A C++ function object passed to the `parallel_for` is called a "SYCL kernel".
A SYCL kernel defines the entry point to the "device part" of the code. The
compiler will emit all symbols accessible from a "kernel". In this code
example, the compiler will emit "foo" function. More details about the
compilation of functions for the device part can be found in the SYCL 1.2.1
specification Section 6.4.
To show to the compiler entry point to the "device part" of the code, the SYCL
runtime can use the `sycl_kernel` attribute in the following way:
```c++
namespace cl {
namespace sycl {
class handler {
template <typename KernelName, typename KernelType/*, ...*/>
__attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
// ...
KernelFuncObj();
}
template <typename KernelName, typename KernelType, int Dims>
void parallel_for(range<Dims> NumWorkItems, KernelType KernelFunc) {
#ifdef __SYCL_DEVICE_ONLY__
sycl_kernel_function<KernelName, KernelType, Dims>(KernelFunc);
#else
// Host implementation
#endif
}
};
} // namespace sycl
} // namespace cl
```
The compiler will also generate an OpenCL kernel using the function marked with
the `sycl_kernel` attribute.
Here is the list of SYCL device compiler expectations with regard to the
function marked with the `sycl_kernel` attribute:
- The function must be a template with at least two type template parameters.
The compiler generates an OpenCL kernel and uses the first template parameter
as a unique name for the generated OpenCL kernel. The host application uses
this unique name to invoke the OpenCL kernel generated for the SYCL kernel
specialized by this name and second template parameter `KernelType` (which
might be an unnamed function object type).
- The function must have at least one parameter. The first parameter is
required to be a function object type (named or unnamed i.e. lambda). The
compiler uses function object type fields to generate OpenCL kernel
parameters.
- The function must return void. The compiler reuses the body of marked functions to
generate the OpenCL kernel body, and the OpenCL kernel must return `void`.
The SYCL kernel in the previous code sample meets these expectations.
}];
}
def DeviceKernelDocs : Documentation {
let Category = DocCatFunction;
let Heading = "device_kernel, nvptx_kernel, amdgpu_kernel, "
"kernel, __kernel";
let Content = [{
These attributes specify that the function represents a kernel for device offloading.
The specific semantics depend on the offloading language, target, and attribute spelling.
Here is a code example using the attribute to mark a function as a kernel:
```c++
[[clang::device_kernel]] int foo(int x) { return ++x; }
```
}];
}
def SYCLExternalDocs : Documentation {
let Category = DocCatFunction;
let Heading = "sycl_external";
let Content = [{
The `sycl_external` attribute indicates that a function defined in another
translation unit may be called by a device function defined in the current
translation unit or, if defined in the current translation unit, the function
may be called by device functions defined in other translation units.
The attribute is intended for use in the implementation of the `SYCL_EXTERNAL`
macro as specified in section 5.10.1, "SYCL functions and member functions
linkage", of the SYCL 2020 specification.
The attribute only appertains to functions and only those that meet the
following requirements:
- Has external linkage
- Is not explicitly defined as deleted (the function may be an explicitly
defaulted function that is defined as deleted)
The attribute shall be present on the first declaration of a function and
may optionally be present on subsequent declarations.
When compiling for a SYCL device target that does not support the generic
address space, the function shall not specify a raw pointer or reference type
as the return type or as a parameter type.
See section 5.10, "SYCL offline linking", of the SYCL 2020 specification.
The following examples demonstrate the use of this attribute:
```c++
[[clang::sycl_external]] void Foo(); // Ok.
[[clang::sycl_external]] void Bar() { /* ... */ } // Ok.
[[clang::sycl_external]] extern void Baz(); // Ok.
[[clang::sycl_external]] static void Quux() { /* ... */ } // error: Quux() has internal linkage.
```
}];
}
def SYCLKernelEntryPointDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
The `sycl_kernel_entry_point` attribute facilitates the launch of a SYCL
kernel and the generation of an offload kernel entry point, sometimes called
a SYCL kernel caller function, suitable for invoking a SYCL kernel on an
offload device. The attribute is intended for use in the implementation of
SYCL kernel invocation functions like the `single_task` and `parallel_for`
member functions of the `sycl::handler` class specified in section 4.9.4,
"Command group `handler` class", of the SYCL 2020 specification.
The attribute requires a single type argument that meets the requirements for
a SYCL kernel name as described in section 5.2, "Naming of kernels", of the
SYCL 2020 specification. A unique kernel name type is required for each
function declared with the attribute. The attribute may not first appear on a
declaration that follows a definition of the function.
The attribute only appertains to functions and only those that meet the
following requirements.
- Has a non-deduced `void` return type.
- Is not a constructor or destructor.
- Is not a non-static member function with an explicit object parameter.
- Is not a C variadic function.
- Is not a coroutine.
- Is not defined as deleted or as defaulted.
- Is not defined with a function try block.
- Is not declared with the `constexpr` or `consteval` specifiers.
- Is not declared with the `[[noreturn]]` attribute.
Use in the implementation of a SYCL kernel invocation function might look as
follows.
```c++
namespace sycl {
class handler {
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char* kernelSymbol, Ts&&... kernelArgs) {
// This code will run on the host and is responsible for calling functions
// appropriate for the desired offload backend (OpenCL, CUDA, HIP,
// Level Zero, etc...) to copy the kernel arguments denoted by kernelArgs
// to a device and to schedule an invocation of the offload kernel entry
// point denoted by kernelSymbol with the copied arguments.
}
template<typename KernelName, typename KernelType>
[[ clang::sycl_kernel_entry_point(KernelName) ]]
void kernel_entry_point(KernelType kernelFunc) {
// This code will run on the device. The call to kernelFunc() invokes
// the SYCL kernel.
kernelFunc();
}
public:
template<typename KernelName, typename KernelType>
void single_task(const KernelType& kernelFunc) {
// This code will run on the host. kernel_entry_point() is called to
// trigger generation of an offload kernel entry point and to schedule
// an invocation of it on a device with kernelFunc (a SYCL kernel object)
// passed as a kernel argument. This call will result in an implicit call
// to sycl_kernel_launch() with the symbol name for the generated offload
// kernel entry point passed as the first function argument followed by
// kernelFunc.
kernel_entry_point<KernelName>(kernelFunc);
}
};
} // namespace sycl
```
A SYCL kernel object is a callable object of class type that is constructed on
a host, often via a lambda expression, and then passed to a SYCL kernel
invocation function to be executed on an offload device. The `kernelFunc`
parameters in the example code above correspond to SYCL kernel objects.
A SYCL kernel object type is required to satisfy the device copyability
requirements specified in section 3.13.1, "Device copyable", of the SYCL 2020
specification. Additionally, any data members of the kernel object type are
required to satisfy section 4.12.4, "Rules for parameter passing to kernels".
For most types, these rules require that the type is trivially copyable.
However, the SYCL specification mandates that certain special SYCL types, such
as `sycl::accessor` and `sycl::stream`, be device copyable even if they are
not trivially copyable. These types require special handling because they cannot
necessarily be copied to device memory as if by `memcpy()`.
The SYCL kernel object and its data members constitute the parameters of an
offload kernel. An offload kernel consists of an offload entry point function
and the set of all functions and variables that are directly or indirectly used
by the entry point function.
A SYCL kernel invocation function is responsible for performing the following
tasks (likely with the help of an offload backend like OpenCL):
1. Identifying the offload kernel entry point to be used for the SYCL kernel.
2. Validating that the SYCL kernel object type and its data members meet the
SYCL device copyability and kernel parameter requirements noted above.
3. Copying the SYCL kernel object and any other kernel arguments to device
memory including any special handling required for SYCL special types.
4. Initiating execution of the offload kernel entry point.
The offload kernel entry point for a SYCL kernel performs the following tasks:
1. Calling the `operator()` member function of the SYCL kernel object.
The `sycl_kernel_entry_point` attribute facilitates or automates these tasks
by providing generation of an offload kernel entry point with a unique symbol
name, type checking of kernel argument requirements, and initiation of kernel
execution via synthesized calls to a `sycl_kernel_launch` template.
A function declared with the `sycl_kernel_entry_point` attribute specifies
the parameters and body of an offload entry point function. Consider the
following call to the `single_task()` SYCL kernel invocation function assuming
an implementation similar to the one shown above.
```c++
struct S { int i; };
void f(sycl::handler &handler, sycl::stream &sout, S s) {
handler.single_task<struct KN>([=] {
sout << "The value of s.i is " << s.i << "\n";
});
}
```
The SYCL kernel object is the result of the lambda expression. The call to
`kernel_entry_point()` via the call to `single_task()` triggers the
generation of an offload kernel entry point function that looks approximately
as follows.
```c++
void sycl-kernel-caller-for-KN(kernel-type kernelFunc) {
kernelFunc();
}
```
There are a few items worthy of note:
1. `sycl-kernel-caller-for-KN` is an exposition only name; the actual name
generated for an entry point is an implementation detail and subject to
change. However, the name will incorporate the SYCL kernel name, `KN`,
that was passed as the `KernelName` template parameter to
`single_task()` and eventually provided as the argument to the
`sycl_kernel_entry_point` attribute in order to ensure that a unique
name is generated for each entry point. There is a one-to-one correspondence
between SYCL kernel names and offload kernel entry points.
2. The SYCL kernel is a lambda closure type and therefore has no name;
`kernel-type` is substituted above and corresponds to the `KernelType`
template parameter deduced in the call to `single_task()`.
3. The parameter and the call to `kernelFunc()` in the function body
correspond to the definition of `kernel_entry_point()` as called by
`single_task()`.
4. The parameter is type checked for conformance with the SYCL device
copyability and kernel parameter requirements.
Within `single_task()`, the call to `kernel_entry_point()` is effectively
replaced with a synthesized call to a ''sycl_kernel_launch\`\` template that
looks approximately as follows.
```c++
sycl_kernel_launch<KN>("sycl-kernel-caller-for-KN", kernelFunc);
```
There are a few items worthy of note:
1. Lookup for the `sycl_kernel_launch` template is performed as if from the
body of the (possibly instantiated) definition of `kernel_entry_point()`.
If name lookup or overload resolution fails, the program is ill-formed.
If the selected overload is a non-static member function, then `this` is
passed as the implicit object parameter.
2. Function arguments passed to `sycl_kernel_launch()` are passed
as if by `std::move(x)`.
3. The `sycl_kernel_launch` template is expected to be provided by the SYCL
library implementation. It is responsible for copying the kernel arguments
to device memory and for scheduling execution of the generated offload
kernel entry point identified by the symbol name passed as the first
function argument. `sycl-kernel-caller-for-KN` is substituted above for
the actual symbol name that would be generated for the offload kernel entry
point.
It is not necessary for a function declared with the `sycl_kernel_entry_point`
attribute to be called for the offload kernel entry point to be emitted. For
inline functions and function templates, any ODR-use will suffice. For other
functions, an ODR-use is not required; the offload kernel entry point will be
emitted if the function is defined. In any case, a call to the function is
required for the synthesized call to `sycl_kernel_launch()` to occur.
A function declared with the `sycl_kernel_entry_point` attribute may include
an exception specification. If a non-throwing exception specification is
present, an exception propagating from the implicit call to the
`sycl_kernel_launch` template will result in a call to `std::terminate()`.
Otherwise, such an exception will propagate normally.
Functions declared with the `sycl_kernel_entry_point` attribute are not
limited to the simple example shown above. They may have additional template
parameters, declare additional function parameters, and have complex control
flow in the function body. The function must abide by the language feature
restrictions described in section 5.4, "Language restrictions for device
functions" in the SYCL 2020 specification. If the function is a non-static
member function, `this` shall not be used in a potentially evaluated
expression.
}];
}
def SYCLSpecialClassDocs : Documentation {
let Category = DocCatStmt;
let Content = [{
SYCL defines some special classes (accessor, sampler, and stream) which require
specific handling during the generation of the SPIR entry point.
The `__attribute__((sycl_special_class))` attribute is used in SYCL
headers to indicate that a class or a struct needs a specific handling when
it is passed from host to device.
Special classes will have a mandatory `__init` method and an optional
`__finalize` method (the `__finalize` method is used only with the
`stream` type). Kernel parameters types are extract from the `__init` method
parameters. The kernel function arguments list is derived from the
arguments of the `__init` method. The arguments of the `__init` method are
copied into the kernel function argument list and the `__init` and
`__finalize` methods are called at the beginning and the end of the kernel,
respectively.
The `__init` and `__finalize` methods must be defined inside the
special class.
Please note that this is an attribute that is used as an internal
implementation detail and not intended to be used by external users.
The syntax of the attribute is as follows:
```text
class __attribute__((sycl_special_class)) accessor {};
class [[clang::sycl_special_class]] accessor {};
```
This is a code example that illustrates the use of the attribute:
```c++
class __attribute__((sycl_special_class)) SpecialType {
int F1;
int F2;
void __init(int f1) {
F1 = f1;
F2 = f1;
}
void __finalize() {}
public:
SpecialType() = default;
int getF2() const { return F2; }
};
int main () {
SpecialType T;
cgh.single_task([=] {
T.getF2();
});
}
```
This would trigger the following kernel entry point in the AST:
```c++
void __sycl_kernel(int f1) {
SpecialType T;
T.__init(f1);
...
T.__finalize()
}
```
}];
}
def C11NoReturnDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
A function declared as `_Noreturn` shall not return to its caller. The
compiler will generate a diagnostic for a function declared as `_Noreturn`
that appears to be capable of returning to its caller. Despite being a type
specifier, the `_Noreturn` attribute cannot be specified on a function
pointer type.
}];
}
def CXX11NoReturnDocs : Documentation {
let Category = DocCatFunction;
let Heading = "noreturn, _Noreturn";
let Content = [{
A function declared as `[[noreturn]]` shall not return to its caller. The
compiler will generate a diagnostic for a function declared as `[[noreturn]]`
that appears to be capable of returning to its caller.
The `[[_Noreturn]]` spelling is deprecated and only exists to ease code
migration for code using `[[noreturn]]` after including `<stdnoreturn.h>`.
}];
}
def NoMergeDocs : Documentation {
let Category = DocCatStmt;
let Content = [{
If a statement is marked `nomerge` and contains call expressions, those call
expressions inside the statement will not be merged during optimization. This
attribute can be used to prevent the optimizer from obscuring the source
location of certain calls. For example, it will prevent tail merging otherwise
identical code sequences that raise an exception or terminate the program. Tail
merging normally reduces the precision of source location information, making
stack traces less useful for debugging. This attribute gives the user control
over the tradeoff between code size and debug information precision.
`nomerge` attribute can also be used as function attribute to prevent all
calls to the specified function from merging. It has no effect on indirect
calls to such functions. For example:
```c++
[[clang::nomerge]] void foo(int) {}
void bar(int x) {
auto *ptr = foo;
if (x) foo(1); else foo(2); // will not be merged
if (x) ptr(1); else ptr(2); // indirect call, can be merged
}
```
`nomerge` attribute can also be used for pointers to functions to
prevent calls through such pointer from merging. In such case the
effect applies only to a specific function pointer. For example:
```c++
[[clang::nomerge]] void (*foo)(int);
void bar(int x) {
auto *ptr = foo;
if (x) foo(1); else foo(2); // will not be merged
if (x) ptr(1); else ptr(2); // 'ptr' has no 'nomerge' attribute, can be merged
}
```
}];
}
def NoInlineDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
This function attribute suppresses the inlining of a function at the call sites
of the function.
`[[clang::noinline]]` spelling can be used as a statement attribute; other
spellings of the attribute are not supported on statements. If a statement is
marked `[[clang::noinline]]` and contains calls, those calls inside the
statement will not be inlined by the compiler.
`__noinline__` can be used as a keyword in CUDA/HIP languages. This is to
avoid diagnostics due to usage of `__attribute__((__noinline__))`
with `__noinline__` defined as a macro as `__attribute__((noinline))`.
```c
int example(void) {
int r;
[[clang::noinline]] foo();
[[clang::noinline]] r = bar();
return r;
}
```