Skip to content

Commit 6b6eee4

Browse files
v-klochkovsarnex
andauthored
[ESIMD][NFC][DOC] Describe the new memory API accepting properties (#13071)
The adds a new file `sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd_functions.md` describing the new memory API including: * block_load(), slm_block_load() * block_store(), slm_block_store() * gather(), slm_gather() * scatter(), slm_scatter() * atomic_update(), slm_atomic_update() * prefetch() Only the block_load/store functions are described in more details at this moment. They include the functions restrictions and hardware requirements/implications from using various features (mask, cache-hints, sizes, etc). --------- Signed-off-by: Klochkov, Vyacheslav N <[email protected]> Co-authored-by: Nick Sarnie <[email protected]>
1 parent 70cad4e commit 6b6eee4

File tree

3 files changed

+688
-15
lines changed

3 files changed

+688
-15
lines changed

sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd.md

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -365,8 +365,8 @@ memory access interface. It supports main SYCL's device memory representations:
365365

366366
Only small subset of `sycl::accessor` APIs is supported in ESIMD context:
367367
- accessor::accessor();
368-
- accessor::get_pointer(); // Supported only with the `-fsycl-esimd-force-stateless-mem` switch.
369-
- accessor::operator[]; // Supported only with the `-fsycl-esimd-force-stateless-mem` switch.
368+
- accessor::get_pointer(); // Supported only with the `-fsycl-esimd-force-stateless-mem` switch (turned ON by default).
369+
- accessor::operator[]; // Supported only with the `-fsycl-esimd-force-stateless-mem` switch (turned ON by default).
370370

371371
ESIMD provides special APIs to access memory through accessors. Those APIs
372372
accept an accessor object as a base reference to the addressed memory and
@@ -419,6 +419,8 @@ They go through extra layer of faster cache.
419419
load/store scalar values through accessors. In case of USM pointers, usual
420420
C++ dereference operator can be used. SLM versions are also available.
421421
422+
See a more detailed list of available memory APIs [here](./sycl_ext_intel_esimd_functions.md).
423+
422424
423425
#### Shared local memory access
424426
@@ -428,13 +430,13 @@ This memory is shared between work items in a workgroup - basically
428430
it is ESIMD variant of the SYCL `local` memory.
429431
430432
SLM variants of APIs have 'slm_' prefix in their names,
431-
e.g. ext::intel::esimd::slm_block_load() or ext::intel::experimental::esimd::lsc_slm_gather().
433+
e.g. ext::intel::esimd::slm_block_load() or ext::intel::esimd::slm_gather().
432434
433435
SLM memory must be explicitly allocated before it is read or written.
434436
435437
There are 3 different ways of SLM allocation in ESIMD:
436-
* static allocation using slm_init<SLMByteSize>() and slm_init(SpecializationConstSLMByteSize)
437-
* semi-dynamic allocation using slm_allocator<SLMByteSize> class
438+
* static allocation using `slm_init<SLMByteSize>()` and `slm_init(SpecializationConstSLMByteSize)`
439+
* semi-dynamic allocation using `slm_allocator<SLMByteSize>` class
438440
* SYCL local accessors
439441
440442
##### Static allocation of SLM using slm_init function.
@@ -457,6 +459,7 @@ Restrictions:
457459
* The call of `slm_init` must be placed in the beginning of the kernel.
458460
If `slm_init` is called in some function 'F' called from kernel, then inlining
459461
of 'F' to the kernel must be forced/guaranteed.
462+
* `slm_init` cannot be used together with `local_accessor` in the same kernel.
460463

461464
##### Semi-dynamic allocation of SLM.
462465
The class `slm_allocator` is designed to be used in basic blocks or functions

0 commit comments

Comments
 (0)